学习型成本模型客户端
地址适用于
libtpu-0.0.40-cp314wheel 中的libtpu.so(BuildID md589edbbe81c5b328a958fe628a9f2207d)。该二进制没有被 strip,下面每个符号都是 demangled C++ 名称。.text/.rodataVMA == 文件偏移;.data.rel.roVMA − 0x200000 == 文件偏移。其他版本会有所不同。
摘要
TPU convolution lowering 路径包含一个完整的、用于 ML 学习型成本模型的客户端侧钩子,但这个构建中没有随附 server、没有 client class,也没有 predictor。该钩子位于 SpatialMajorConvolution::ComputeWindowConfigInternal(@ 0x13172c80)内部,也就是搜索 window-tiling 空间以寻找最快 convolution schedule 的函数。钩子会访问一个借用的 EmitterLearnedCostModelBase* 指针,该指针存储在 emitter 的 this+0x20d0;当该指针非空且 enable_learned_cost_model 字节被设置时,搜索会通过四个虚函数槽位咨询学习型模型:逐 instruction 的启用检查(vtable+0x10)、候选 window 注册调用(vtable+0x18,RegisterCandidateWindow)、最快 window 状态查询(vtable+0x30,其 absl::Status 结果在学习型答案和解析型答案之间选择),以及仅在非错误状态下执行的配对取值(vtable+0x38),后者返回学习得到的周期估计(返回对象 +0x138 处的 float)。该指针一路从 80 参数的 LoweringEmitter ctor 经由 ConvolutionEmitter::Create 借用传下去。
熟悉的参照框架是 XLA 的可插拔成本模型接口:带有虚拟 Predict/Register 方法的抽象基类、把 features 批处理为 RPC 发送给 embedding/inference 服务的具体客户端,以及在任何 RPC 失败时由 absl::Status 门控回退到解析模型。libtpu-0.0.40 随附的是接口契约(proto options、四枚举 mode 状态机、gflag wiring、调用点、注册结果上的 CHECK_OK,以及 Failed to get fastest window using learned cost model 失败日志回退),但没有实现:没有 xla::jellyfish::EmitterLearnedCostModelBase vtable,没有具体的 LearnedCostModelClient,没有 LearnedCostModelService::Stub,也没有嵌入的模型字节。因此在发布二进制中 EmitterLearnedCostModelBase* 始终为 null,所以每次咨询都会短路,并且原样使用解析型 TpuHloCostAnalysis / window-search 结果。
本页记录真实且可重新实现的内容:wire-level options proto(LearnedCostModelClientOptions 及其 ServiceType enum、RPC endpoint 字段、恢复的 C++ struct layout)、接口实现的四枚举 mode/validation/DB-query 状态机、恢复的 EmitterLearnedCostModelBase vtable 形状(四个被使用的槽位)、客户端接收的 MLCostModelWindowInfo 请求 payload、每个调用点的 status/error handling,以及把整个层降回解析模型的精确 null/flag 双门控。
对重新实现而言,契约是:
EmitterLearnedCostModelBasevtable:IsEnabled(HloInstruction*)@+0x10(返回bool)、RegisterCandidateWindow(LcmKey, MLCostModelWindowInfo)@+0x18(返回absl::Status,被CHECK_OK)、最快 window status query @+0x30(返回absl::Status,软失败),以及配对 value-fetch @+0x38(返回指向结果对象的指针,该对象在+0x138处以 float 保存学习得到的周期估计)。MLCostModelWindowInfo请求 payload,也就是客户端为每个候选 window 接收的八个 designated-initializer 字段,包括estimated_cycles_classic(解析型答案,始终作为下限提供)。- status handling:
+0x30查询结果absl::Status选择 learned-vs-analytic;非 OK 时该层记录spatial_major_convolution.cc:4006并带着 classic search result 落入SetupBestConfig。 - 通往解析模型的双门控:
this+0x20d0null check 和TpuCompEnv+0xed6(enable_learned_cost_model)字节;任意一个为 false 都降到 analytic。 LearnedCostModelClientOptionsproto wire shape、ServiceType(LOCAL/REMOTE)enum,以及恢复的 struct layout,也就是客户端实现必须从 gflag 解析的内容。
| 钩子位置 | SpatialMajorConvolution::ComputeWindowConfigInternal @ 0x13172c80 |
| 注册 lambda | …::ComputeWindowConfigInternal(…)::$_0(policy_func)@ 0x1317fe00 |
| 客户端指针槽位 | this+0x20d0(SpatialMajorConvolution),借用的 EmitterLearnedCostModelBase* |
| 启用槽位 | vtable+0x10,bool IsEnabled(HloInstruction*)(@ call 0x13173…) |
| 注册槽位 | vtable+0x18,absl::Status RegisterCandidateWindow(LcmKey, MLCostModelWindowInfo) |
| 查询槽位 | vtable+0x30,absl::Status 最快 window status query(软失败) |
| 取值槽位 | vtable+0x38,学习得到的周期估计的配对 fetch(float @ result +0x138) |
| 启用 flag 字节 | *(byte*)(GetTpuCompEnv(inst)+0xed6) == enable_learned_cost_model |
| Options proto / vtable | xla::jellyfish::LearnedCostModelClientOptions @ 0x21cffc10 |
| gflag | xla_tpu_emitter_learned_cost_model_options = AutoOr<EmitterLearnedCostModelOptions> |
| 失败日志回退 | spatial_major_convolution.cc:4006 → SetupBestConfig(analytic) |
| 发布默认值 | client pointer 处处为 null → 始终使用 analytic model |
随附什么与不随附什么
学习型成本模型是一个典型的“未来扩展钩子”:schema、gflag、consumer call site 和 failure-fallback 都随附;predictor 不随附。这一拆分是精确的,并且可通过符号扫描验证。
| 组件 | 存在? | 状态 |
|---|---|---|
EmitterLearnedCostModelOptions proto | 是 | 可达 |
LearnedCostModelClientOptions proto | 是 | 可达 |
FusionDataProtoGenerationOptions proto | 是 | 可达 |
EmbeddingCacheEntry / EmbeddingCacheDB protos | 是 | 可达 |
LearnedCostModelMode / DbQueryType / MLOutputValidationStrategy / ServiceType enums | 是 | 已解码 |
gflag xla_tpu_emitter_learned_cost_model_options | 是 | 已解析 |
| Consumer call sites(4 个 vtable 槽位) | 是 | 代码存在,运行时死亡 |
RegisterCandidateWindow 结果上的 CHECK_OK | 是 | Source spatial_major_convolution.cc:3996 |
Failed to get fastest window … analytic fallback | 是 | Source spatial_major_convolution.cc:4006 |
EmitterLearnedCostModelBase vtable / typeinfo | 否 | 仅类型 |
具体 LearnedCostModelClient class | 否 | 缺失 |
LearnedCostModelService::Stub(gRPC) | 否 | 缺失 |
Predict / Inference / Score / EstimateCycles method | 否 | 缺失 |
| 嵌入模型(SavedModel / ONNX / TFLite blob) | 否 | 缺失 |
NOTE: 这种缺失是正向证据,不是分析空白。扫描(未 strip 的)符号表可以找到
…ClientOptionsproto 族及其 ser/deser 方法,但EmitterLearnedCostModelBase只作为函数参数类型存在于ConvolutionEmitter::Create、SpatialMajorConvolution::SpatialMajorConvolution和LoweringEmitter::LoweringEmitter的 mangled names 中。一个仅以借用指针使用的 class 不需要在 consumer 的 translation unit 中发出 vtable 或 typeinfo,这正是只有 out-of-tree 实现的接口会留下的 footprint。
客户端接口:EmitterLearnedCostModelBase vtable
目的
EmitterLearnedCostModelBase 是 convolution emitter 调用的抽象接口。它由 emitter 借用(从不拥有),因此没有被使用的 destructor 槽位。反编译代码触达了三个虚函数槽位;其签名由调用点寄存器使用和 CHECK_OK 字面量固定。
恢复的 vtable
| 槽位 | 方法(推断) | 返回值 | 调用点证据 |
|---|---|---|---|
+0x00 | offset_to_top | — | Itanium ABI |
+0x08 | typeinfo ptr | — | Itanium ABI |
+0x10 | bool IsEnabled(const HloInstruction*) | bool(在 %al 中) | (*(…)(*v44 + 16))(v44, hlo),arg 是 *(HloInstruction**)(this+72);结果驱动 consult 分支 |
+0x18 | absl::Status RegisterCandidateWindow(const LcmKey&, const MLCostModelWindowInfo&) | absl::Status | lambda 中的 (*(…)(*v52 + 24))(v52, key.first, key.second, &status);被 CHECK_OK |
+0x30 | fastest-window status query | absl::Status(stack-returned StatusRep*) | (*(…)(**(this+0x20d0) + 48))(&out, this_lcm, fp.first, fp.second);status != OK 选择 fallback(记录 .cc:4006) |
+0x38 | paired value-fetch(学习得到的周期估计) | 指向结果对象的指针 | (*(…)(**(this+0x20d0) + 56))(this_lcm, fp.first, fp.second);仅当 +0x30 OK 时到达,检查 [1] presence,并读取 result +0x138 处的 float |
window search 内的调用序列
// SpatialMajorConvolution::ComputeWindowConfigInternal @0x13172c80
bool consult = false; // v172
void* lcm = *(void**)(this + 0x20d0); // borrowed EmitterLearnedCostModelBase*
if (lcm) { // null → analytic (default ship)
if (enable_learned_cost_model) { // bool param a9, see gate below
const HloInstruction* hlo = *(const HloInstruction**)(this + 72);
consult = (*(bool(**)(void*, const HloInstruction*))(*(void**)lcm + 0x10))(lcm, hlo); // IsEnabled
if (consult) {
fp = xla::GetHloInstructionFingerprint(hlo); // @0x13180b80 — the prediction key
}
} else {
consult = false;
}
}
// ... the window-tiling search runs (IterateThroughWindowConfigs), invoking the
// registration lambda once per candidate window when consult is true ...
if (consult != true) { // learned model not used
SetupBestConfig(/* classic search result */); // analytic answer
} else {
// fastest-window status query @ vtable+0x30, keyed by the fingerprint
status = (*(Status(**)(…))(*(void**)lcm + 0x30))(&out, lcm, fp.first, fp.second);
if (!status.ok()) { // analytic fallback
LOG("Failed to get fastest window using learned cost model "
"for instruction: ", hlo, " with status: ", status); // .cc:4006
SetupBestConfig(/* classic search result */); // <-- fall through to analytic
} else {
// paired value-fetch @ vtable+0x38, same fingerprint key
result = (*(void*(**)(…))(*(void**)lcm + 0x38))(lcm, fp.first, fp.second);
float learned_cycles = (result[1] ? *(float*)((char*)*result + 0x138) : 0.0f);
SetupBestConfig(/* learned-selected window, learned_cycles */);
}
}
```text
> **GOTCHA:** `vtable+0x10` 槽位接受 `HloInstruction*`(从 `this+72` 加载)并返回 `bool`;不要把它误认为无参数 `IsEnabled()`。逐 instruction 参数使真实客户端能够根据 op shape/type 门控学习型路径。在发布二进制中该槽位永远不会到达,因为 `lcm` 为 null。
### 注册调用:`RegisterCandidateWindow`
window-search 迭代器会为每个候选 tiling 调用一次 `std::function`(`…::$_0` policy_func @ `0x1317fe00`)。当学习型路径处于活动状态时,该 lambda 根据其参数在栈上构建一个 `MLCostModelWindowInfo`(InlinedVectors 通过 `inlined_vector_internal::Storage::InitFrom` 深拷贝),并调用 `vtable+0x18`:
```c
// lambda body @0x1317fe00 — one call per candidate window
if (*(byte*)(GetTpuCompEnv(hlo) + 0xed6) == 1 // enable_learned_cost_model
&& capture_flag == 1 && a8 < threshold) { // **(this+0x80)+0x10 byte + window bound
void* lcm = *(void**)(this + 0x20d0);
Status s = (*(Status(**)(void*, long, long, Status*))(*(void**)lcm + 0x18))(
lcm, lcm_key.first, lcm_key.second, &out); // RegisterCandidateWindow
// CHECK_OK — fatal on a non-OK registration:
// "learned_cost_model_->RegisterCandidateWindow( *lcm_key,
// MLCostModelWindowInfo( {.activations_window = …, .kernel_window = …,
// .output_window = …, .iteration_bounds = …, .window_info = …,
// .vmem_footprint_granules = estimated_granules, .bundles = estimated_bundles,
// .estimated_cycles_classic = estimated_cycles})) is OK" // .cc:3996
}NOTE:
RegisterCandidateWindow被CHECK_OK(失败时 fatal),而最快 window 查询 @+0x30是软失败(记录日志并回退)。这种不对称是有意的:注册是向考虑集合中做本地 bookkeeping push(不应失败),而预测查询可能发生瞬态失败(RPC down),并且必须降级到解析模型而不是中止编译。
请求 Payload:MLCostModelWindowInfo
由 CHECK_OK designated-initializer 字面量和 lambda 的栈构建序列重建。这是客户端接收的逐候选 window feature record;尾部的 estimated_cycles_classic 是解析型答案,作为 baseline/floor 提供。
struct MLCostModelWindowInfo {
absl::InlinedVector<int64_t, 6> activations_window; // input window dims (deep-copied)
absl::InlinedVector<int64_t, 6> kernel_window; // kernel window dims
absl::InlinedVector<int64_t, 6> output_window; // output window dims
absl::InlinedVector<int64_t, 6> iteration_bounds; // outer loop bounds
WindowSizingInfo window_info; // sizing metadata (copied from a7)
int64_t vmem_footprint_granules; // estimated_granules
int64_t bundles; // estimated_bundles
int64_t estimated_cycles_classic; // analytic-model estimate (baseline)
};
```text
`RegisterCandidateWindow` 的第一个参数是 `*lcm_key`,也就是一个 `LcmKey`,其两个 8 字节半部(`v53 = *v50; v54 = v50[1]`)按值传递。它的完整形状无法恢复,除了知道它是一个 two-word key;它会在一次 fusion search 内对候选去重,并且很可能编码 operation type、MXU format 和逐 window hash。相反,最快 window 查询 @ `+0x30` 以 `xla::GetHloInstructionFingerprint(hlo)`(@ `0x13180b80`)为键,该键也作为 two-word value(`fp.first`, `fp.second`)传递。
> **CONTRACT:** 解析搜索考虑的每个候选 window 都会连同其 classic cycle estimate 一起报告给客户端。因此真实客户端从不冷启动:它可以原样返回 classic number(validation strategy `ALWAYS_TRUST` 关闭)、对其 clamp(`NO_NEGATIVE_CYCLES`),或覆盖它。这就是下面 `MLOutputValidationStrategy` enum 的 wire-level 含义。
---
## Mode 状态机:四个枚举
options proto 编码了一个由(缺失的)客户端实现的小状态机。所有四个 enum 都从嵌入的 `FileDescriptorProto` 解码;value-name 字符串存在于 `.rodata` 中。
### `LearnedCostModelMode`
| 值 | 名称 | 语义 |
|---|---|---|
| 0 | `LEARNED_COST_MODEL_MODE_INVALID` | 默认,视为“没有学习型成本模型” |
| 1 | `LEARNED_COST_MODEL_MODE_ONLY_DB` | 只从预构建 DB 查找周期 |
| 2 | `LEARNED_COST_MODEL_MODE_ONLY_ML_PREDICTION` | 始终使用 ML predictor |
| 3 | `LEARNED_COST_MODEL_MODE_DB_WITH_FALLBACK_TO_ML_PREDICTION` | DB 优先,未命中时使用 ML |
| 4 | `LEARNED_COST_MODEL_MODE_ONLY_DATA_COLLECTION` | 转储 `FusionData` protos 供离线训练;不评分 |
mode 名称揭示了预期设计:一个离线 DB,保存预先测量的 `(window-config → cycles)` tuple,并由 ML predictor 填补空缺。`+0x18` `RegisterCandidateWindow` 槽位是数据收集 / 考虑集合 push(所有 mode 都使用);`+0x30` 查询槽位是 DB-lookup-or-predict。
### `MLOutputValidationStrategy`
| 值 | 名称 | 语义 |
|---|---|---|
| 0 | `ML_OUTPUT_VALIDATION_STRATEGY_NONE` | 不验证 |
| 1 | `ML_OUTPUT_VALIDATION_STRATEGY_NEVER_TRUST` | 始终回退到 classic cost model |
| 2 | `ML_OUTPUT_VALIDATION_STRATEGY_ALWAYS_TRUST` | 原样采用 ML output |
| 3 | `ML_OUTPUT_VALIDATION_STRATEGY_NO_NEGATIVE_CYCLES` | 只拒绝负周期预测 |
### `DbQueryType`
| 值 | 名称 | 语义 |
|---|---|---|
| 0 | `DB_QUERY_TYPE_NONE` | 不查询 DB |
| 1 | `DB_QUERY_TYPE_REPLAY_PREDICITIONS` |(sic,proto 拼写错误)重放已存储的 ML cycles |
| 2 | `DB_QUERY_TYPE_GROUND_TRUTH` | 查找已测量的 ground-truth cycles |
### `LearnedCostModelClientOptions.ServiceType`
| 值 | 名称 | 语义 |
|---|---|---|
| 0 | `SERVICE_TYPE_UNSPECIFIED` | 无效 sentinel |
| 1 | `SERVICE_TYPE_LOCAL` | 从 `local_embedding_model_path` 加载模型,进程内运行 |
| 2 | `SERVICE_TYPE_REMOTE` | 向 `remote_embedding_server_address` 发起 RPC |
> **GOTCHA:** 这个 wheel 中 `SERVICE_TYPE_LOCAL` 和 `SERVICE_TYPE_REMOTE` 都不可构建。LOCAL 路径没有进程内模型加载器,REMOTE 路径也没有 `LearnedCostModelService::Stub` / `BlockingUnaryCall`。二进制*确实*随附无关的 gRPC stubs(`BarnaCoreInterWorkerCommunicationRpc::Stub`、`RuntimeMetricService::Stub`、`MegaScaleTransport::Stub`),这证明如果存在学习型成本模型 RPC stub,它会可见。
---
## Options Proto Wire Shape
RPC/service 配置完全作为一个 gflag 中的序列化 proto 承载,没有专用 boolean 或 endpoint flag。两个 proto 文件由一个 sub-message 连接:`EmitterLearnedCostModelOptions.learned_cost_model_client_options` 保存一个 `LearnedCostModelClientOptions`。
```protobuf
package xla.jellyfish;
import "third_party/tensorflow/core/framework/tensor.proto";
message LearnedCostModelClientOptions {
enum ServiceType { SERVICE_TYPE_UNSPECIFIED = 0; SERVICE_TYPE_LOCAL = 1; SERVICE_TYPE_REMOTE = 2; }
optional ServiceType embedding_service_type = 1;
optional string remote_embedding_server_address = 2; // REMOTE RPC endpoint
optional string remote_embedding_model_name = 3; // REMOTE model selector
optional int32 inflight_rpc_monitoring_interval_milliseconds = 4; // RPC liveness poll (serialized int32)
optional string local_embedding_model_path = 5; // LOCAL model file
optional string embedding_cache_path = 6; // EmbeddingCacheDB on disk
optional FusionDataProtoGenerationOptions fusion_data_proto_generation_options = 7;
optional int32 max_batch_size = 8; // RPC batch size (serialized int32)
}
message EmbeddingCacheEntry { optional bytes fingerprint = 1; optional tensorflow.TensorProto embedding = 2; }
message EmbeddingCacheDB { repeated EmbeddingCacheEntry entries = 1; }
message FusionDataProtoGenerationOptions {
optional bool include_standalone_fusion_module = 1;
optional bool include_expert_and_gating_features = 2;
}C++ struct layout,从 copy-ctor LearnedCostModelClientOptions(Arena*, const&) @ 0x1db653e0 字节精确恢复(存储的 vtable 是 off_21CFFC20):
| 偏移 | 字段 | 大小 | 备注 |
|---|---|---|---|
+0x00 | vtable ptr | 8 | 0x21cffc10+0x10 |
+0x08 | internal::InternalMetadata | 8 | arena / unknown-field tag |
+0x10 | uint32_t _has_bits_ | 4 | presence bitmap |
+0x14 | cached_size | 4 | |
+0x18 | TaggedStringPtr remote_embedding_server_address | 8 | ForceCopy if tag bits set |
+0x20 | TaggedStringPtr remote_embedding_model_name | 8 | |
+0x28 | TaggedStringPtr local_embedding_model_path | 8 | |
+0x30 | TaggedStringPtr embedding_cache_path | 8 | |
+0x38 | FusionDataProtoGenerationOptions* | 8 | copied iff _has_bits_ & 0x10 |
+0x40 | int32_t embedding_service_type(enum,field 1) | 4 | serialized first,tag byte 0x08,from *((int*)this+16) |
+0x44 | int32_t inflight_rpc_monitoring_interval_milliseconds(field 4) | 4 | WriteInt32ToArrayWithField<4> from *((int*)this+17) |
+0x48 | int32_t max_batch_size(field 8) | 4 | WriteInt32ToArrayWithField<8> from *((int*)this+18) |
NOTE: copy-ctor 用
(*(byte*)(this+0x10) & 0x10)保护 sub-message,_has_bits_bit 4 控制fusion_data_proto_generation_options。四个 string 字段是proto2::internal::TaggedStringPtr,并在其低 tag bits 被设置时被ForceCopy(arena-owned vs inline)。三个尾部整数由_InternalSerialize@0x1db65920固定:field 1(embedding_service_type)从+0x40写入,field 4(inflight_rpc_monitoring_interval_milliseconds)从+0x44写入,field 8(max_batch_size)从+0x48写入,全部通过WriteInt32ToArrayWithField,因此两个 interval/batch 字段被序列化为 32 位,尽管下面的.proto文本声明它们为int64。copy-ctor 把+0x40/+0x44这一对作为一个 8 字节 word 复制,并把+0x48作为单独 dword 复制。两个非 enum 整数的 proto 声明宽度为中等置信度(serializer 使用 int32 路径;对小值而言 on-wire varint 与宽度无关)。
拥有它的 EmitterLearnedCostModelOptions 添加了顶层开关:enable_learned_cost_model(tag 1,gate byte)、cost_model_mode、db_query_type、ml_output_validation_strategy、db_path、max_num_considered_windows、dump_fusion_data_proto[_dir]。其中只有 enable_learned_cost_model 有 runtime consumer(+0xed6 gate);其余字段会被反序列化,但由于它们配置的客户端缺失而死亡。
解析型回退:双门控
两个独立 gate 会把整个学习型层降回解析型 window search。任意一个为 false 就足够。
// Gate 1 — pointer null check (ComputeWindowConfigInternal @0x13172c80)
void* lcm = *(void**)(this + 0x20d0);
if (!lcm) consult = false; // DEFAULT in shipping libtpu-0.0.40 — always taken
// Gate 2 — enable flag (lambda @0x1317fe00, and the IsEnabled branch)
if (*(byte*)(GetTpuCompEnv(hlo) + 0xed6) != 1) /* skip */ ; // enable_learned_cost_model
// Gate 3 (defence-in-depth) — captured byte flag inside the lambda
if (**(byte**)(capture + 0x80 ... + 0x10) != 1) /* skip */ ;
```text
`this+0x20d0` 指针逐字从 `EmitterLearnedCostModelBase*` 构造函数参数设置(`SpatialMajorConvolution` C2 @ `0x130dd180` 用原始 `mov` 存储它,没有分配),该参数由 `ConvolutionEmitter::Create`(@ `0x130d86c0`)传递 null,并最终来自 `LoweringEmitter::LoweringEmitter`(@ `0x10c309c0`)。由于这个构建中没有调用方曾提供非 null 指针,Gate 1 始终触发,而 Gate 2–3 不可达:解析型 [TpuHloCostAnalysis](tpu-hlo-cost-analysis.md) flop 模型和 classic window search(`SetupBestConfig`)驱动每个 convolution-lowering 决策。
当存在客户端并且查询失败时,`spatial_major_convolution.cc:4006` 处的软回退会记录 instruction 和失败的 `absl::Status`,随后用 classic search result 调用 `SetupBestConfig`,也就是 Gate 1 到达的同一代码路径。因此学习型模型 RPC outage 在功能上等同于完全没有客户端。
| Gate | 位置 | 字段 | 发布值 | 为 false 时的效果 |
|---|---|---|---|---|
| 1 — pointer | `this+0x20d0` | 借用的 `EmitterLearnedCostModelBase*` | **null** | 跳过 consult;analytic |
| 2 — enable | `GetTpuCompEnv(hlo)+0xed6` | `enable_learned_cost_model` | 0(proto 默认值) | 跳过 register/query |
| 3 — capture | lambda capture `+0x10` byte | propagated enable | 0 | 跳过逐候选 register |
| soft — query | `+0x30` result `absl::Status` | n/a(不可达) | OK | 记录 `.cc:4006`,analytic |
---
## 函数与符号映射
| 符号 | 地址 | 角色 |
|---|---|---|
| `SpatialMajorConvolution::ComputeWindowConfigInternal` | `0x13172c80` | hook site:enable check、fingerprint、query、fallback |
| `…::ComputeWindowConfigInternal(…)::$_0`(policy_func) | `0x1317fe00` | 逐候选 `RegisterCandidateWindow` 调用 + `CHECK_OK` |
| `SpatialMajorConvolution::SpatialMajorConvolution`(C2) | `0x130dd180` | 在 `this+0x20d0` 存储 `EmitterLearnedCostModelBase*` |
| `ConvolutionEmitter::Create` | `0x130d86c0` | 转发(null)client pointer |
| `LoweringEmitter::LoweringEmitter`(C1) | `0x10c309c0` | 产生借用的 client pointer |
| `xla::GetHloInstructionFingerprint` | `0x13180b80` | 构建 `+0x30` query key |
| `LearnedCostModelClientOptions(Arena*, const&)` | `0x1db653e0` | copy-ctor → struct layout |
| `LearnedCostModelClientOptions::_InternalSerialize` | `0x1db65920` | proto wire encode |
| `EmitterLearnedCostModelOptions(Arena*)` | `0x1db63f20` | owning options proto ctor |
| `AutoOr<EmitterLearnedCostModelOptions>::ParseFlag` | `0x1d745680` | gflag → proto parse |
| `LearnedCostModelClientOptions` vtable | `0x21cffc10` | proto vtable |
| `EmitterLearnedCostModelBase` vtable / typeinfo | — | **不存在**(仅接口) |
| `LearnedCostModelClient` concrete class | — | **不存在** |
| `LearnedCostModelService::Stub`(gRPC) | — | **不存在** |
> **QUIRK:** 失败日志调用点在这个构建(`0.0.40`)中显示为 `spatial_major_convolution.cc:4006`,`RegisterCandidateWindow` `CHECK_OK` 显示为 `:3996`。源代码行号随构建版本而变;周围 VA 和 wire contract 才是稳定锚点。
---
## 交叉引用
- [成本模型概览](overview.md):为什么随附模型是数据表驱动而非 ML;本页是该声明背后的细节
- [TpuHloCostAnalysis](tpu-hlo-cost-analysis.md):回退使用的解析型 flop/byte 模型,以及 `estimated_cycles_classic` 来源
- [归一化计算成本](normalized-computation-cost.md):classic window search 供给的 convolution-cycle cache
- [成本模型日志](cost-model-logging.md):兄弟 `AutoOr<…Options>` gflag consumer pattern(impure logging options)
- [逐 opcode 周期常量](per-opcode-cycle-constants.md):解析估计所引用的逐世代周期表