Skip to content

PhaseCompile 扩展(类型 9)

本页中的所有地址均适用于 libtpu-0.0.40-cp314 wheel(libtpu_lts_20260413_b_RC00,BuildID 89edbbe81c5b328a958fe628a9f2207d,ELF x86-64 DYN,约 745 MB)里的 libtpu.so。PJRT C-API 表面是 v0.103。其他构建会重新编号槽位并偏移地址。

摘要

PhaseCompile 扩展PJRT_Extension_Type == 9,一个 64 字节结构体)是 libtpu 的增量式编译 ABI。一次性 PJRT_Client_Compile(槽位 25)会把整个程序和一个 CompileOptionsProto 交给客户端,并返回完成的、已加载到设备的可执行体;而 PhaseCompile 扩展允许宿主框架一次驱动一个具名阶段来运行同一个 xla::TpuCompiler,例如先运行 phase0_stablehlo_to_hlo,缓存中间 HLO blob,之后再运行 phase1_hlo_opts 以及后续阶段。阶段边界是序列化的 PjRtPartialProgramProto blob,因此部分结果可以被缓存、检查,或拆分到不同宿主上执行。两个表面都由同一个 xla::TpuCompiler 对象和同一个 CompileOptionsProto 反序列化关口支撑;它们的差异只在于中间产物是否暴露出来。

该扩展是一个扁平的函数指针表,附加在 PJRT_Api 扩展链上。五个槽位中有三个是标准通用 XLA C-ABI 驱动(pjrt::PJRT_PhaseCompile_*,逐字节等同于上游 pjrt_c_api_phase_compile_internal.cc 代码);另两个 Get_CompilerDestroy_Compiler 由 TPU 注入(tpu_plugin::*TpuPhaseCompiler),也是该表面上唯一的 TPU 专用代码。注入的 Get_Compiler 会创建一个 xla::TpuCompiler,并在其上注册恰好六个具名阶段;随后通用的 Run_Phase / Get_Phase_Names 驱动通过已注册编译器上的两个虚槽(+48 GetPhaseNames+56 RunPhases)分发。

本页负责说明PhaseCompile 扩展结构体、五槽位阶段 API,以及编译选项 proto 摄入路径,也就是每一次 PhaseCompile 调用与整体式路径共享的部分。编译生成的最终可执行体,包括它的 C-ABI 盒子、serialize / deserialize 往返以及 Execute,见可执行体加载与执行。这些阶段实际驱动的 HLO pass 管线(phase1_hlo_opts 的主体和 lowering 阶段)见 HLO Pass 注册表;缓存部分结果和完整结果的内容寻址存储见编译缓存

对于重新实现,契约是:

  • 64 字节扩展结构体。 一个 PJRT_Extension_Base 头(size=64type=9next),后跟五个函数指针,位于 +0x18 … +0x38。创建器从参数写入两个 TPU 注入的指针,并硬编码三个通用驱动。
  • 编译器句柄。 Get_Compiler 通过 operator new 分配一个 0x40 字节的 xla::TpuCompiler(vtable off_2177D660),调用 RegisterAllPhases,把它包装进一个 0x10 字节的 { base, owner } holder,然后把 holder 存入 args 输出字段。宿主持有的不透明句柄是 holder,而不是编译器;Run_Phase/Get_Phase_Names 解引用 holder.baseDestroy_Compiler 释放 holder.owner
  • 六个已注册阶段、它们的顺序,以及 phase3_linking / phase3_linking_test_only 是同一个函数的两次注册,仅绑定的 bool 不同。
  • 共享的编译选项关口。 Run_Phase(和每个编译入口一样)执行 CompileOptionsProto::ParseFromStringCompileOptions::FromProto,解析失败时返回完全相同的 "PJRT_Client_Compile: failed to deserialize CompileOptionsProto"
  • char-buffer 编组。 阶段名和部分程序 blob 以并行的 (char**, size_t*, count) span 跨越 C-ABI;Convert* 辅助函数在它们和 C++ string / PjRtPartialProgramProto vector 之间转换,C_Buffers_Destroy 释放返回的数组。
扩展类型9 (PJRT_Extension_Type::PhaseCompile)
扩展结构体大小0x40 (64 B); 5 个 fn-ptr 位于 +0x18..+0x38
创建器pjrt::CreatePhaseCompileExtension @ 0x0E6F42A0
存储位置.bss slot GetTpuPjrtApi::phase_compile_extension @ 0x224C3B18
Get_Compiler (+0x18, TPU)pjrt::tpu_plugin::GetTpuPhaseCompiler @ 0x0E6AA320
Destroy_Compiler (+0x20, TPU)pjrt::tpu_plugin::DestroyTpuPhaseCompiler @ 0x0E6AA400
Run_Phase (+0x28, generic)pjrt::PJRT_PhaseCompile_Run_Phase @ 0x0E6F42E0
Get_Phase_Names (+0x30, generic)pjrt::PJRT_PhaseCompile_Get_Phase_Names @ 0x0E6F4A60
C_Buffers_Destroy (+0x38, generic)pjrt::PJRT_PhaseCompile_C_Buffers_Destroy @ 0x0E6F4CC0
编译器对象xla::TpuCompiler (0x40 B, vtable off_2177D660)
阶段数量6(5 个不同函数 + 1 个 test_only 别名)
来源(断言)…/pjrt/c/pjrt_c_api_phase_compile_internal.cc; …/research/pjrt/tpu_pjrt_compiler.cc

扩展结构体

目的

PhaseCompile 扩展是 PJRT_Api 扩展链中的一个链节(见扩展链):宿主遍历 PJRT_Extension_Basenext 列表,匹配 type == 9,然后把该链节转换为 PJRT_PhaseCompile_Extension 以访问五个阶段函数。该结构体在插件初始化时由 CreatePhaseCompileExtension 组装一次,放入固定的 .bss 槽位(0x224C3B18),并由 GetTpuPjrtApi 使用作为参数传入的两个 TPU 注入函数指针进行填充。

布局

c
struct PJRT_PhaseCompile_Extension {       // struct_size 0x40 (64 B), type 9
    PJRT_Extension_Base base;              // +0x00 size=64; +0x08 type=9; +0x10 next
    /* +0x18 */ Get_Compiler*      get_compiler;       // TPU-injected (arg a3)
    /* +0x20 */ Destroy_Compiler*  destroy_compiler;   // TPU-injected (arg a4)
    /* +0x28 */ Run_Phase*         run_phase;          // generic pjrt:: driver
    /* +0x30 */ Get_Phase_Names*   get_phase_names;    // generic pjrt:: driver
    /* +0x38 */ C_Buffers_Destroy* c_buffers_destroy;  // generic pjrt:: driver
};

算法 — 创建器

c
function CreatePhaseCompileExtension(ext, next, get_fn, destroy_fn):  // 0xe6f42a0
    ext[+0x00] = 64                                  // struct_size
    ext[+0x08] = 9                                   // PJRT_Extension_Type::PhaseCompile
    ext[+0x10] = next                                // chain link (arg a2)
    ext[+0x18] = get_fn                              // arg a3  (TPU GetTpuPhaseCompiler)
    ext[+0x20] = destroy_fn                          // arg a4  (TPU DestroyTpuPhaseCompiler)
    ext[+0x28] = pjrt::PJRT_PhaseCompile_Run_Phase           // generic driver
    ext[+0x30] = pjrt::PJRT_PhaseCompile_Get_Phase_Names     // generic driver
    ext[+0x38] = pjrt::PJRT_PhaseCompile_C_Buffers_Destroy   // generic driver
    return ext

创建器是纯粹的扁平表写入器,没有分配也没有分支;在 0x0E6F42A0 可验证为八次 store 和一次 return。注入get/destroy)和通用run/names/destroy-buffers)之间的拆分就是整个设计:后端只提供编译器工厂和销毁逻辑,上游 C-ABI 驱动负责针对工厂返回的任何编译器对象完成所有编组。

注意 — 两个注入指针以调用参数(a3a4)形式到达,而不是来自创建器直接引用的 TPU 符号。这是 libtpu 扩展链中共享的通用创建器模式:CreatePhaseCompileExtension 本身是标准 XLA 代码,TPU GetTpuPjrtApi 会用 TPU 函数指针调用它。将 libtpu 移植到另一个后端的重新实现者只需重写 Get_Compiler/Destroy_Compiler;创建器和三个驱动保持不变。

函数映射

槽位字段符号地址来源
+0x18get_compilertpu_plugin::GetTpuPhaseCompiler0x0E6AA320TPU 注入
+0x20destroy_compilertpu_plugin::DestroyTpuPhaseCompiler0x0E6AA400TPU 注入
+0x28run_phasepjrt::PJRT_PhaseCompile_Run_Phase0x0E6F42E0通用
+0x30get_phase_namespjrt::PJRT_PhaseCompile_Get_Phase_Names0x0E6F4A60通用
+0x38c_buffers_destroypjrt::PJRT_PhaseCompile_C_Buffers_Destroy0x0E6F4CC0通用
创建器pjrt::CreatePhaseCompileExtension0x0E6F42A0通用

Get_Compiler — 创建阶段编译器

目的

Get_Compiler 是宿主首先调用的入口:它生成每个后续 Run_Phase / Get_Phase_Names 调用都会携带的不透明编译器句柄。对于 TPU 后端,这是 tpu_plugin::GetTpuPhaseCompiler0x0E6AA320),它构建一个新的 xla::TpuCompiler,在其上注册六个阶段,并把它装箱到一个很小的所有权 holder 中。

入口点

text
GetTpuPhaseCompiler (0xe6aa320, ext +0x18)
  ── ActualStructSizeIsGreaterOrEqual("PJRT_PhaseCompile_Get_Compiler_Args", 35, 24)
  ── operator new(0x40)                  ── xla::TpuCompiler, vtable off_2177D660, body zeroed
  ── xla::TpuCompiler::RegisterAllPhases (0xf849ec0)   ── registers the 6 phases
  └─ operator new(0x10)                  ── holder { base, owner } → args[+0x10]

算法

c
function GetTpuPhaseCompiler(args):                       // 0xe6aa320
    if !ActualStructSizeIsGreaterOrEqual(                 // line 14
            "PJRT_PhaseCompile_Get_Compiler_Args", min=35, cur=24, args->struct_size):
        return new PJRT_Error{status}                     // guard miss

    compiler = operator new(0x40)                         // xla::TpuCompiler
    zero(compiler+0x08 .. compiler+0x3F)                  // vmovups ymm 32B x2
    compiler->vtable = off_2177D660                       // line 24

    status = xla::TpuCompiler::RegisterAllPhases(compiler) // 0xf849ec0, line 26
    if status == OK:                                      // (== 1)
        holder = operator new(0x10)                       // line 29
        holder->base  = compiler                          //   +0x00 — Run/Names deref this
        holder->owner = compiler                          //   +0x08 — Destroy frees this
        args[+0x10]   = holder                            // args[2], line 32
        return NULL                                       // success
    else:
        err = new PJRT_Error{status}                      // line 38
        compiler->vtable[+8](compiler)                    // TpuCompiler dtor — tear down
        return err

注意事项

  • 0x10 holder 是句柄,不是编译器。 宿主接收的是 holder,即 { TpuCompiler* base; TpuCompiler* owner; } 这一层间接引用。Run_Phase/Get_Phase_Names 读取 holder->base 并通过它的 vtable 分发;Destroy_Compiler 置空并释放 holder->owner。创建时两个指针相同,但 holder 的存在使所有权(要删除什么)与访问权(要调用什么)可以分离,这是上游 PJRT 用于可转移但可借用句柄的模式。
  • 失败路径会干净销毁。 如果 RegisterAllPhases 失败,半构造的编译器会在返回错误前通过 vtable+8(其 dtor)销毁,不会泄漏这个 0x40 字节对象。错误是堆上的 PJRT_Error,持有 absl::Status,形状与每个 PJRT wrapper 返回的错误相同。
  • struct_size 包络。 args 结构体 guard 是 min=35, cur=24。与该表面其他位置一样(见 API Vtable 重建),这是 ActualStructSizeIsGreaterOrEqual(name, expected_size_for_min_version, current_size, host_supplied_size);结构体小于最小值的宿主会被 PJRT_Error 拒绝,主体不会读取任何内容。

Destroy_Compiler

c
function DestroyTpuPhaseCompiler(args):                   // 0xe6aa400
    holder = args[+0x10]                                  // args[2]
    if holder:
        owner = holder->owner                             // +0x08
        holder->owner = NULL
        if owner: owner->vtable[+8](owner)                // TpuCompiler dtor
        free(holder)                                      // free the 0x10 box

Destroy_Compiler 没有 struct_size guard,返回 void,它是纯清理例程。它通过 Get_Compiler 在失败路径上使用的同一个 vtable+8 槽位释放编译器,然后释放 holder。在 dtor 调用前置空 owner,可以防止宿主错误地销毁两次时出现 double-free。


Run_Phase — 驱动一个具名阶段

目的

Run_Phasepjrt::PJRT_PhaseCompile_Run_Phase @ 0x0E6F42E0)是通用驱动,用来针对编译器句柄执行一个有序的具名阶段列表。它是增量 ABI 的核心:它从 C-ABI char-buffer span 反编组输入的部分程序 blob 和阶段名列表,反序列化 CompileOptionsProto(共享关口),分发到 RunPhases(编译器 vtable+56),再把输出的部分程序 vector 重新编组回 char-buffer span。

入口点

text
PJRT_PhaseCompile_Run_Phase (0xe6f42e0, ext +0x28)
  ── ActualStructSizeIsGreaterOrEqual("PJRT_PhaseCompile_Run_Phase_Args", 32, 120)
  ── ConvertCharBuffersToCppStrings           (phase names    → vector<string>)
  ── ConvertCharBuffersToPjRtPartialProgramProtos (input progs → vector<PjRtPartialProgramProto>)
  ── CompileOptionsProto::ParseFromString → CompileOptions::FromProto   [shared chokepoint]
  ── null-check holder (args[+0x10])          ── "…Run_Phase: phase compiler is null"
  ── (*holder->base->vtable[+56])(…)          ── PjRtPhaseCompiler::RunPhases
  └─ ConvertPjRtPartialProgramProtosToCharBuffers (output vec → char buffers, written out)

算法

c
function PJRT_PhaseCompile_Run_Phase(args):               // 0xe6f42e0
    if !ActualStructSizeIsGreaterOrEqual(                 // line 76
            "PJRT_PhaseCompile_Run_Phase_Args", min=32, cur=120, args->struct_size):
        return new PJRT_Error{status}

    // --- unmarshal the two parallel char-buffer spans ---
    phase_names = ConvertCharBuffersToCppStrings(         // line 84
                      ptrs=args[+0x30], count=args[+0x40],
                      sizes=args[+0x38])                  // → vector<string>
    input_programs = ConvertCharBuffersToPjRtPartialProgramProtos(  // line 85
                      ptrs=args[+0x18], count=args[+0x28],
                      sizes=args[+0x20])                  // → vector<PjRtPartialProgramProto>

    // --- the compile-options chokepoint (see §Compile-Options Ingest) ---
    proto = CompileOptionsProto{}                          // line 134
    if !proto.ParseFromString(args[+0x48], args[+0x50]):   // line 135 — ptr +0x48, len +0x50
        return new PJRT_Error{MakeErrorImpl<3>(
            "PJRT_Client_Compile: failed to deserialize CompileOptionsProto",
            …, "…/pjrt_c_api_phase_compile_internal.cc")}   // :46 (NOT wrapper_impl:1113)
    opts = CompileOptions::FromProto(proto)                // line 138

    holder = args[+0x10]                                   // args[2]
    if !holder:                                            // line ~230
        return new PJRT_Error{MakeErrorImpl<13>(
            "PJRT_PhaseCompile_Run_Phase: phase compiler is null",  // :77
            …, "…/pjrt_c_api_phase_compile_internal.cc")}

    // --- drive the pipeline: RunPhases via compiler vtable +56 ---
    out_programs = (*holder->base->vtable[+56])(           // line 238
                       holder->base, &opts, input_programs,
                       topology = *(args[+0x88]) + 8,      //   PjRtTopologyDescription
                       phase_names)                        // → StatusOr<vector<PjRtPartialProgramProto>>
    if !out_programs.ok(): return new PJRT_Error{out_programs.status}

    // --- remarshal results into the out char-buffer span ---
    ConvertPjRtPartialProgramProtosToCharBuffers(out_programs, args+0x60..)
    return NULL

args 布局

反编译出的主体会读取 PJRT_PhaseCompile_Run_Phase_Args 内的下列偏移。两个并行 char-buffer span(程序和阶段名)各自携带一个 (ptrs, sizes, count) 三元组;这两个 span 的语义标签是从 Convert* 调用参数顺序推断出来的。

偏移字段(恢复)含义
+0x00struct_sizeguard 输入;min 32,cur 120
+0x10phase_compiler来自 Get_Compiler{base,owner} holder(args[2]
+0x18input_programschar** — 序列化的 PjRtPartialProgramProto blob
+0x20input_programs_sizessize_t* — 每个 blob 的字节长度
+0x28input_programs_count元素数量
+0x30phase_nameschar** — 阶段名字符串
+0x38phase_names_sizessize_t* — 每个名称的长度
+0x40phase_names_count元素数量
+0x60out_programschar** out — 成功时写入
+0x68out_programs_sizessize_t* out
+0x70out_programs_countcount out
+0x88topology*(args+0x88)+8PjRtTopologyDescription
+0x48compile_optionschar* — 序列化的 CompileOptionsProto,送入 ParseFromString
+0x50compile_options_sizesize_t — proto 字节长度

注意 — 上面的偏移直接来自 Run_Phase 主体 0xe6f42e0(第 84-85、134-135 行):两个 char-buffer span 是不同的 (ptr, size, count) 三元组,程序位于 (+0x18, +0x20, +0x28),阶段名位于 (+0x30, +0x38, +0x40);选项 proto 是 (ptr=+0x48, len=+0x50)。反编译器把这些渲染为十进制(*(a1+72)/*(a1+80) = 十六进制 +0x48/+0x50);不要把十进制误读为十六进制。确切的逐字段名称取决于公开的 pjrt_c_api_phase_compile.h 头文件顺序,而它不在二进制中;偏移具有高置信度。

注意事项

  • 分发是 vtable+56 反编译主体调用 (*(holder->base->vtable + 56))(…),即 PjRtPhaseCompiler::RunPhases @ 0x1D16AF20。这是整体式 PJRT_Client_Compile 驱动的同一个编译器对象,只是通过不同 vtable 槽位到达;分阶段 ABI 暴露整体式路径会丢弃的中间 PjRtPartialProgramProto 产物。
  • 错误字符串有误导性,它即使在这里也命名为 PJRT_Client_Compile CompileOptionsProto 解析失败会返回 "PJRT_Client_Compile: failed to deserialize CompileOptionsProto",虽然这里是 Run_Phase。解析逻辑是每个编译入口都会运行的同一个双调用序列,但 Run_Phase 携带其自己的内联副本(发射在 pjrt_c_api_phase_compile_internal.cc:46MakeErrorImpl<3>),不同于整体式入口共享的 pjrt_c_api_wrapper_impl.cc:1113 位置;这个字符串只是逐字复制(和 DeserializeAndLoad 的怪癖相同,见可执行体执行)。重新实现者不应基于该字符串进行错误处理。
  • 每一步的错误都会以 PJRT_Error* 暴露。 反编组失败、选项解析失败、编译器为空,或 RunPhases 失败,都会生成一个堆上的 PJRT_Error,其中持有 absl::Status;成功返回 NULL

Get_Phase_Names — 枚举已注册阶段

目的

Get_Phase_Namespjrt::PJRT_PhaseCompile_Get_Phase_Names @ 0x0E6F4A60)返回编译器注册的有序阶段名列表,让宿主可以在驱动管线前发现它。它是 Run_Phase 的只读配套入口。

算法

c
function PJRT_PhaseCompile_Get_Phase_Names(args):         // 0xe6f4a60
    if !ActualStructSizeIsGreaterOrEqual(                 // line 25
            "PJRT_PhaseCompile_Get_Phase_Names_Args", min=38, cur=48, args->struct_size):
        return new PJRT_Error{status}

    holder = args[+0x10]                                   // args[2]
    if !holder:                                            // line 34
        return new PJRT_Error{MakeErrorImpl<13>(
            "PJRT_PhaseCompile_Get_Phase_Names: phase compiler is null",  // :100
            …, "…/pjrt_c_api_phase_compile_internal.cc")}

    names = (*holder->base->vtable[+48])(holder->base)     // PjRtPhaseCompiler::GetPhaseNames
                                                           //   0x1d16ae20, line 40
    args[+0x18] = ConvertCppStringsToCharBuffers(names, &args[+0x20])  // ptrs / sizes
    args[+0x28] = names.count                              // line 71
    free(names)                                            // line 60-114: release the vector
    return NULL

注意事项

  • 分发是 vtable+48 GetPhaseNames @ 0x1D16AE20 返回已注册阶段名 vector(即 RegisterAllPhases 追加它们的顺序)。随后驱动把它们通过 ConvertCppStringsToCharBuffers 转为 args[+0x18] / args[+0x20](指针数组 / sizes),并把 count 放在 args[+0x28]
  • args 结构体名称有两种拼写。 struct_size 断言字符串是 "PJRT_PhaseCompile_Get_Phase_Names_Args"0xe6f4a60,第 25 行),但 IDA 恢复出的函数 mangled symbol 接收 PJRT_PhaseCompile_Get_PhaseNames_Args*PhaseNames 之间没有下划线)。两种拼写都出现在二进制中;线上结构体使用断言命名的带下划线形式。
  • 返回的 vector 会就地释放。 转换成 char buffer 后,主体会遍历 vector<string>,释放每个长字符串分配(libc++ SSO 符号位检查 (char)x < 0),再释放 backing array,也就是标准的 0xAAAAAAAAAAAAAAAB reciprocal-of-24 stride math。宿主只拥有转换后的 char buffer,之后通过 C_Buffers_Destroy 释放它们。

C_Buffers_Destroy — 释放返回的数组

目的

C_Buffers_Destroypjrt::PJRT_PhaseCompile_C_Buffers_Destroy @ 0x0E6F4CC0)释放 Run_Phaseout_programs)和 Get_Phase_Names(阶段名)返回给宿主的 char-buffer 数组。它是宿主针对这些分配必须调用的清理钩子。

算法

c
function PJRT_PhaseCompile_C_Buffers_Destroy(args):       // 0xe6f4cc0
    ptrs  = args[+0x10]                                   // args[2] — char** array
    count = args[+0x20]                                   // args[4] — element count
    if ptrs:
        for i in 0 .. count-1: free(ptrs[i])              // each buffer
        free(ptrs)                                        // the array
    sizes = args[+0x18]                                   // args[3] — size_t* array
    if sizes: free(sizes)
    // no status returned — pure cleanup

它不返回状态;与 Destroy_Compiler 一样,它是 void 清理例程。重新实现必须为每个产生 buffer 的 Run_Phase / Get_Phase_Names 配对调用一次 C_Buffers_Destroy,否则会泄漏返回的数组。


六个已注册阶段

目的

xla::TpuCompiler::RegisterAllPhases @ 0x0F849EC0 是 TPU 后端声明其分阶段管线的位置。它调用 PjRtPhaseCompiler::RegisterPhase @ 0x1D16AB20 六次,每次传入一个名称、一个编译函数,以及共享的 CommonPhaseValidatorRegisterPhase 会把 {name → CompilationPhaseFunctions{compile_fn, validator_fn}} 插入 hash map,并把名称追加到 GetPhaseNames 返回的有序 vector 中。

阶段表

这六个注册在 RegisterAllPhases0xf849ec0,六次 RegisterPhase 调用)中得到验证;名称字符串在 .rodata 中逐字存在。

顺序阶段名编译函数地址输出格式
1phase0_stablehlo_to_hloCompilePhase0StablehloToHlo0x0F84DE60unopt_hlo
2phase1_hlo_optsCompilePhase1HloOptimizations0x0F84EE00optimized HLO
3phase2a_tlp_loweringCompilePhase2aTlpLowering0x0F850840TLP-lowered
4phase2b_deduped_loweringCompilePhase2bDedupedLowering0x0F852180deduped-lowered
5phase3_linkingCompilePhase3Linking(test=false)0x0F852F40device program
6phase3_linking_test_onlyCompilePhase3Linking(test=true)0x0F852F40device program

算法 — 注册

c
function TpuCompiler::RegisterAllPhases(this):            // 0xf849ec0
    RegisterPhase(this, "phase0_stablehlo_to_hlo",        // line 74-80
                  CompilePhase0StablehloToHlo,  CommonPhaseValidator)
    RegisterPhase(this, "phase1_hlo_opts",                // line 101-108
                  CompilePhase1HloOptimizations, CommonPhaseValidator)
    RegisterPhase(this, "phase2a_tlp_lowering",           // line 137-143
                  CompilePhase2aTlpLowering,    CommonPhaseValidator)
    RegisterPhase(this, "phase2b_deduped_lowering",       // line 175-181
                  CompilePhase2bDedupedLowering, CommonPhaseValidator)
    RegisterPhase(this, "phase3_linking",                 // line 204-212
                  bind_front(CompilePhase3Linking, false), CommonPhaseValidator)
    RegisterPhase(this, "phase3_linking_test_only",       // line 237-244
                  bind_front(CompilePhase3Linking, true),  CommonPhaseValidator)
    // each returns Status; a failure aborts and propagates (Get_Compiler tears down)

每个阶段函数具有如下签名

c
StatusOr<vector<PjRtPartialProgramProto>>
    CompilePhaseN(CompileOptions,
                  Span<const PjRtPartialProgramProto> /* input */,
                  const PjRtTopologyDescription&);

共享 validator 是 Status CommonPhaseValidator(Span<const PjRtPartialProgramProto>),在每个阶段前针对输入程序运行。RegisterPhase 将两者存入一个 CompilationPhaseFunctions 记录;RunPhases 在该记录上通过 vtable+72 调用 validator,并通过 vtable+40 调用编译函数。

怪癖 — phase3_linkingphase3_linking_test_only同一个函数CompilePhase3Linking @ 0xf852f40)的两次注册,仅绑定的 bool test 参数不同(false 生产,true 测试)。它们不是五阶段管线加第六阶段;它们是五阶段管线(phase0phase1phase2aphase2bphase3)以及终端 linker 的一个测试别名。通过 Get_Phase_Names 枚举阶段的宿主会看到六个名称;正常编译驱动五个。

驱动 — RunPhases

c
function PjRtPhaseCompiler::RunPhases(opts, programs, topology, phase_names):  // 0x1d16af20
    current = programs                                    // vector<PjRtPartialProgramProto>
    for name in phase_names:                              // host-supplied order
        rec = phase_map.find(name)
        if !rec:
            return Error("No phase compiler/validator registered with phase name \"%s\"", name)  // :281
        rec->validator(current)                           // CompilationPhaseFunctions vtable +72
        current = rec->compile_fn(opts, current, topology) // vtable +40 → new vector
    return current                                        // final artifacts

每个阶段的输出会成为下一阶段的输入(vector 被重新赋值),因此宿主可以单独运行 ["phase0_stablehlo_to_hlo"],持久化 unopt_hlo blob,之后再从该 blob 恢复并运行 ["phase1_hlo_opts", "phase2a_tlp_lowering", …]RegisterPhase 会用 "A phase compiler/validator with Phase name \"%s\" already exists"pjrt_compiler.cc:259)拒绝重复名称,并用第 250/253/256 行的 guard 拒绝空名称或空函数。

部分程序产物格式

跨越每个阶段边界的 blob 是一个 PjRtPartialProgramProto(proto3,package xlapjrt_partial_program.proto):

text
message PjRtPartialProgramProto {
  bytes  program          = 1;   // serialized stage output (e.g. unopt HLO)
  string program_format   = 2;   // e.g. "unopt_hlo"
  string producer_phase   = 3;   // e.g. "phase0_stablehlo_to_hlo"
  repeated string consumer_phases = 4;  // e.g. ["phase1_hlo_opts"]
  string version          = 5;
  string program_name     = 6;
}

CompilePhase0StablehloToHlo0xf84de60)是一个完整示例:它对输入 StableHLO module 执行 ParseMlirModuleString,使用 getDefaultChloToHighLevelMhloOptions 通过 MlirToXlaComputation 生成 HloModuleProto,把 arg/output layout-mode 和 memory-space 属性填入 HLO 的 frontend_attributes(键为 kArgLayoutModesAttr / kOutLayoutModesAttr / kArgMemorySpacesAttr / kOutMemorySpacesAttr),并发出 PjRtPartialProgramProto{ program=serialized HLO, program_format="unopt_hlo", producer_phase="phase0_stablehlo_to_hlo", consumer_phases=["phase1_hlo_opts"] }。下游阶段的 HLO-pass 内容见 HLO Pass 注册表

注意 — phase1_hlo_optsphase2a_tlp_loweringphase2b_deduped_loweringphase3_linking 的内部主体在此只按输入/输出契约层面刻画(每个都把 vector<PjRtPartialProgramProto> 映射到 vector<PjRtPartialProgramProto>);已确认 phase 3 会构建 BufferAssignmentProto 并发出可加载的设备程序。它们的内部 pass 管线很大,位于编译器页面,而不在这里。


编译选项摄入

目的

libtpu 中的每个编译入口,包括 PJRT_Client_CompilePJRT_CompilePJRT_Client_LoadPJRT_Executable_DeserializeAndLoad 以及 PhaseCompile Run_Phase,都会通过同一个双调用序列CompileOptionsProto::ParseFromStringCompileOptions::FromProto)反序列化同一个 CompileOptionsProto,并发出逐字节相同的失败字符串。前四个共享一个发射位置(pjrt_c_api_wrapper_impl.cc:1113);Run_Phase 用自己的位置(pjrt_c_api_phase_compile_internal.cc:46)内联重新实现该序列,逻辑相同、诊断相同、源码行不同。这是线上选项变成活的 C++ xla::CompileOptions 的汇聚点。TPU 编译器配置(拥有 1,121 个字段的 xla.jellyfish.TpuCompilationEnvironment)以及 xla::DebugOptions / XLA flags 全部都位于这个 proto 内部。

关口

c
// PJRT_Client_Compile / PJRT_Compile / PJRT_Client_Load / DeserializeAndLoad → wrapper_impl.cc:1113
// PhaseCompile Run_Phase (line 134) → its own inline copy at phase_compile_internal.cc:46
proto = CompileOptionsProto{}                          // arena = 0
if !proto2::MessageLite::ParseFromString(&proto, ptr, len):
    return new PJRT_Error{MakeErrorImpl<3>(
        "PJRT_Client_Compile: failed to deserialize CompileOptionsProto",  // verbatim, all 5
        …, /* wrapper_impl.cc:1113  | Run_Phase: phase_compile_internal.cc:46 */)}
opts = xla::CompileOptions::FromProto(proto)           // wire → live C++ options

线格式 schema

CompileOptionsProto(proto3,package xla…/pjrt/proto/compile_options.proto)是载体。TPU 编译环境在其中嵌套三层:

text
CompileOptionsProto {
  repeated ShapeProto argument_layouts        = 1;
  bool   parameter_is_tupled_arguments        = 2;
  ExecutableBuildOptionsProto executable_build_options = 3;  // ← env carrier
  bool   compile_portable_executable          = 4;
  int64  profile_version                      = 5;
  bytes  serialized_multi_slice_config        = 6;
  repeated EnvOptionOverridesEntry env_option_overrides = 7; // ← per-flag overrides
  GpuTargetConfigProto target_config          = 8;
  bool   allow_in_place_mlir_modification     = 9;
  PrecisionConfig.Precision matrix_unit_operand_precision = 10;
  string compiler_variant                     = 11;
}
ExecutableBuildOptionsProto {
  int64  device_ordinal                       = 1;
  DebugOptions debug_options                  = 3;   // ← XLA-standard flags
  CompilationEnvironmentsProto comp_envs      = 13;  // ← THE env nest
  int64  num_replicas = 4; num_partitions = 5;
  bool   use_spmd_partitioning = 6; use_auto_spmd_partitioning = 7;
  bytes  fdo_profile                          = 14;  // feedback-directed-opt
  ...
}
CompilationEnvironmentsProto { repeated google.protobuf.Any environments = 1; }

三种注入通道

重新实现者需要知道,TPU 编译器配置到达编译器有三种同心方式,从最窄到最宽:

通道位置作用域
单 flag 覆盖CompileOptionsProto.env_option_overrides(字段 7)一个 flag,本次编译;map<name, {string|bool|int|double}>
完整 env,按编译executable_build_options.comp_envs.environments[Any → TpuCompilationEnvironment](3→13→1)整个 1,121 字段表,本次编译
进程全局TpuExecutable ext SetTpuCompilationEnv(类型 17,槽位 +0x400xe6dd400一个 CompilationEnvironments singleton,之后的编译会继承

前两个位于 CompileOptionsProto 内部,因此和其他每个编译入口一样,以完全相同的方式通过 Run_Phase 的关口。第三个是 TpuExecutable 扩展(类型 17)上的独立带外表面,不是 PhaseCompile 扩展;它在那个扩展处记录,这里只作提示,让重新实现者知道 PhaseCompile 路径并不是唯一的 env 通道。

嵌套环境的解码

FromProto 运行时,CompilationEnvironments::CreateFromProto @ 0x1E63E5A0 会遍历 comp_envs.environments 中的每个 Any,通过 FindMessageTypeByName 在生成的 descriptor pool 中解析其 type_url(例如 xla.jellyfish.TpuCompilationEnvironment),用 MessageFactory::New 创建 prototype,对 Any.value 字节执行 InternalUnpackTo,并用 AddEnvImpl 把活的 message 加入 FlatHashMap<Descriptor*, unique_ptr<Message>>。格式错误的 Any 会产生 compilation_environments.cc:139..160 处的诊断("Invalid/Unknown/Unsupported CompilationEnvironment message type""Unable to unpack…")。因此 TpuCompilationEnvironment 是在编译时通过反射解码的,即按 protobuf 字段名查询,而不是按固定偏移查询。完整字段表见编译器页面。

怪癖 — libtpu 的编译完全在进程内进行:CompileOptionsProto 在同一地址空间中被消费并通过反射解码;没有编译器子进程。xla::CommonPjRtClient::supports_two_phase_launch @ 0xe6edbc0 硬返回 1,但 "two-phase" 指的是编译/加载是可分离的 API 调用,而不是编译会 fork 到进程外。(这与某些其他 PJRT 后端相反,那些后端的 Compile 会 fork 一个外部编译器二进制,并以 argv/临时文件传递选项。libtpu 的 TPU 编译器是静态链接的。)


与整体式路径的关系

PhaseCompile 扩展和一次性 PJRT_Client_Compile(槽位 25)是同一个编译器之上的两个 ABI:

方面整体式(PJRT_Client_Compile,槽位 25)分阶段(PhaseCompile ext,类型 9)
调用者提供整个程序 + CompileOptionsProto逐阶段 PjRtPartialProgramProto blob + 阶段名列表 + 选项
返回完成的、已加载到设备的 PJRT_LoadedExecutable阶段之间的中间 PjRtPartialProgramProto 产物
底层对象xla::TpuCompiler(通过 client vtable)xla::TpuCompiler(通过 ext Get_Compiler holder)
编译器分发client Compile 虚函数在已注册阶段上执行 RunPhasesvtable+56
中间产物丢弃暴露(可缓存、可检查、可跨宿主拆分)
选项关口CompileOptionsProto::ParseFromStringFromProto相同

分阶段 ABI 提供了整体式 ABI 无法提供的三件事:(a)缓存部分编译的程序(在 phase0 后持久化 unopt_hlo blob),(b)跨宿主拆分阶段,以及(c)当只有后期阶段选项改变时,只重新编译下游阶段。任一路径生成的最终可执行体,包括其 serialize/deserialize 往返和 Execute,见可执行体加载与执行


相关组件

名称关系
xla::TpuCompilerGet_Compiler 创建并由 RegisterAllPhases 配置的 0x40 字节编译器对象
xla::PjRtPhaseCompiler提供 RegisterPhase / RunPhases / GetPhaseNames 的基类(vtable +40/+48/+56/+72
CompilationPhaseFunctions按阶段名存储的 {compile_fn, validator_fn} 记录
pjrt::ActualStructSizeIsGreaterOrEqual每个入口的 struct_size 版本 guard,与每个 PJRT wrapper 共享
xla::CompilationEnvironmentsFromProto 时解码嵌套的 TpuCompilationEnvironment Any
TpuExecutable 扩展(类型 17)另一个 env 通道(SetTpuCompilationEnv),相对于编译是带外的

交叉引用

  • 扩展链 — 宿主如何遍历 PJRT_Extension_Base.next 并匹配 type == 9 以到达该结构体
  • API Vtable 重建PJRT_Api 槽位表,以及每个入口共享的 struct_size 向后兼容 guard
  • 可执行体加载与执行 — 编译生成的最终可执行体:PJRT_Client_Compile、serialize / deserialize 往返,以及 Execute
  • 客户端与设备PJRT_Client_Create,构建 TpuClient 并安装 env 解码所用 descriptor pool 的注入槽位
  • PJRT 概览 — 编译表面位于插件生命周期中的位置
  • HLO Pass 注册表phase1/phase2/phase3 主体驱动的 HLO pass 管线
  • 编译缓存 — 缓存部分(PjRtPartialProgramProto)和完整结果的内容寻址存储