动态形状支持

下面的符号名、VA 和 build-id 适用于来自 libtpu-0.0.40-cp314 wheel 的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）。其他版本会不同；请把每个 VA 都视为版本固定。

摘要

TPU 不会分配字节大小依赖运行时值的缓冲区。它的整个动态形状策略是静态填充到边界，并在定宽元数据前缀中携带真实运行时大小。标记为动态的维度携带一个静态上界（存储在 xla::Shape 中的 extent）以及一个按维度的 is_dynamic 位；运行时值位于 [0, bound]。编译器按上界确定每个缓冲区的大小，而 xla::DynamicPadder pass 会在 HLO 预 pass 集合中运行两次，重写模块，使其运行之后，布局分配、tiling、MSA 和 LLO lowering 都看到一个完全静态的模块。动态性只以两种形式保留下来：(a) 计算每个维度运行时大小的普通 S32 标量 SSA 值，以及 (b) 两个边界 custom-call，PadToStatic 和 SliceToDynamic，它们读写预置在每个动态缓冲区前面的 1024 字节元数据前缀。

本页记录重新实现时必须正确处理的三层内容：有界动态 Shape 模型和大小操作数穿线（DynamicParameterBinding、DynamicDimensionInference、SetDimensionSize/GetDimensionSize）；DynamicPadder pass 及其配套项（DynamicIndexSplitter、DynamicDimensionSimplifier）——它们重写什么，以及在哪里插入边界 custom-call；以及填充到静态 tile 的策略——前缀如何定大小，Target::ShapeWithMetadataSizeBytes 如何计算缓冲区物理字节数，以及两个边界 emitter 如何把前缀 lowering 到 LLO。window/conv same-padding 的边界算术和 SparseCore MLIR 动态维度子系统不在本页范围内，仅在分支处注明。

重新实现必须遵守的契约：

静态填充到上界，绝不进行运行时定大小分配。 对动态 shape，Target::ShapeWithMetadataSizeBytes(shape) 是 ShapeSizeBytesRaw(shape) + prefix；对静态 shape，则是 ShapeSizeBytesRaw(shape)。ShapeSizeBytesRaw 使用静态（上界）extent。没有任何代码路径会让缓冲区的物理大小成为运行时标量的函数。
元数据前缀是编译期常量：1024 字节。 Target::DynamicShapeMetadataPrefixBytes() 在每一代 TPU 上都返回 1024，并断言可被 sizeof(int32_t) 整除（256 个 int32 维度大小槽）且 <= ChunkSizeBytes()。
DynamicPadder 是主 lowering pass；它在布局之前运行。 它位于 PreOptimization 阶段的第 21 步（在第 20 步的 DynamicDimensionSimplifier 之后，第 4 步还有更早的 DynamicIndexSplitter — 见 hlo-pre-passes.md）。它之后只剩下 Set/GetDimensionSize S32 标量和两个边界 custom-call。
维度大小是穿过 HLO 图的 S32 标量。 入口参数动态维度通过 DynamicParameterBinding 绑定到另一个 S32 入口参数；图内由 DynamicDimensionInference map 携带每个维度的大小 SSA 值；一个有 37 个 override 的前向 visitor 传播它们。
TPU 拒绝无界动态性。 只有有界动态维度（带静态上界的动态维度）可编译；unbounded dynamism is not supported 是门禁。


主 pass	`xla::DynamicPadder::RunImpl(HloModule*, exec_threads)` @ `0x16998ca0`（约 5,918 行反编译代码）
配套项 — 第 4 步	`xla::DynamicIndexSplitter::RunImpl` @ `0x164ae740`
配套项 — 第 20 步	`xla::DynamicDimensionSimplifier::RunImpl` @ `0x164d0020`
维度大小分析	`xla::DynamicDimensionInference::Run` @ `0x1e39ad20`；visitor `DynamicDimensionInferenceVisitor::Run` @ `0x1e3984c0`
边界插入 helper	`DynamicDimensionInferenceVisitor::RequiresPadToStatic` @ `0x1e39a7e0`；`::InsertPadToStaticOnInstruction` @ `0x1e390920`
去填充重写器	`(anon)::DynamicShapeRemovingVisitor::ConvertToDynamic` @ `0x169a7bc0`
边界 emitter (LLO)	`jellyfish::PadToStaticEmitter::Emit` @ `0x10c9ad40`；`jellyfish::SliceToDynamicEmitter::Emit` @ `0x10c9c6c0`
定大小	`Target::ShapeWithMetadataSizeBytes` @ `0x1d619f20` → `TransferSizeUtil::ShapeWithMetadataSizeBytes` @ `0x1d6aea00`；`Target::DynamicShapeMetadataPrefixBytes` @ `0x1d61c4e0`；`Target::DynamicShapeMetadataPrefixShape` @ `0x1d61c500`
源码单元	`third_party/tensorflow/compiler/xla/service/dynamic_padder.cc`（字符串锚定）
置信度	CONFIRMED（字节锚定），除非某行或 callout 另有说明

有界动态 Shape 模型

该模型是标准 XLA 按维度动态位与作为上界的静态 extent 配对。每个 xla::Shape 都对每个维度 i 携带一个静态 extent D[i] 和一个布尔值 is_dynamic[i]。当 is_dynamic[i] 为 true 时，D[i] 是上界，运行时值位于 [0, D[i]]。相关访问器都有字节锚点：

方法	VA	作用
`xla::Shape::is_dynamic_dimension(int)`	`0x1e52c9e0`	查询按维度动态位
`xla::Shape::set_dynamic_dimension(int, bool)`	`0x20cd8c60`	设置按维度动态位
`xla::Shape::is_static()`	`0x20cd8f80`	当且仅当没有维度是动态时为 true

动态维度的运行时大小始终是一个 S32 标量。两个维度大小 op 的 shape-inference helper 固定了这一点：

xla::ShapeInference::InferSetDimensionSizeShape(Shape, Shape, long) @ 0x1e541d20
xla::ShapeInference::InferGetDimensionSizeShape(Shape, long) @ 0x1e541960

诊断 SetDimensionSize's value has to be S32 scalar, got %s 强制操作数 dtype。动态维度 op 家族的 HloOpcode 字节如下（来自 opcode 目录交叉引用的权威映射）：

text

0x35 dynamic-reshape       0x36 dynamic-slice        0x37 dynamic-update-slice
0x3f get-dimension-size    0x70 set-dimension-size   0x51 pad
```text

> **NOTE — 反编译交叉检查。** `DynamicPadder::RunImpl` (`0x16998ca0`) 通过原始 opcode 区分这些 op：在约第 1077 行测试 `v25 == 63 || v25 == 112`（十进制 `0x3f`/`0x70` = get-/set-dimension-size），并在判断一条指令是否“已经静态”之前特殊处理 `SetBound` custom-call。[Confidence: CONFIRMED — 直接从反编译 switch 读取。]

上界上的编译期不变量由逐字诊断守卫（重写器绝不会静默截断）：

- `Dimension size has to be less-equal than upper bound %lld for dimension %lld in shape %s`
- `dynamic size must be less than or equal to static size`
- `Shape size has to be less than %d in dynamic shape bounded by %s`
- `Shape size has to be greater or equal than 0 in dynamic shape bounded by %s`
- `Non-positive constant for dynamic size`
- `requires 'shape' to have at most one dynamic dimension, but got multiple dynamic dimensions at indices {0} and {1}` — 某些 op 最多允许一个动态维度。

jellyfish 层还把按维度的 tiling/并行 extent 建模为 `std::variant<long, jellyfish::DynamicBound>` — tiling 维度要么是具体的 `long` extent，要么是 `DynamicBound`（有界动态 extent）。这出现在 `VerifyParallelAttributes(..., absl::Span<variant<long, DynamicBound> const>, ...)` (`0x14516f40`) 以及用于上界下 literal 求值的 `LiteralBase::Piece::CopyElementsWithDynamicBound<T>` 中。[Confidence: HIGH — 符号锚定，函数体未完全追踪。]

---

## 维度大小穿线

动态维度大小作为普通 S32 标量 SSA 值穿过三层。

### 第 1 层 — 入口绑定：`DynamicParameterBinding`

对于维度为动态的入口参数，运行时大小由*另一个入口参数*（S32 标量）提供。绑定是 `map<DynamicDimension{param, index, dim} -> DynamicSizeParameter{param, index}>`。逐字诊断锚定了该机制：

```text
-- Input param number %lld at %s has dim %lld as dynamic dimension,
   which is represented by param number %lld at %s

DynamicParameterBinding 是分析种子：DynamicDimensionInferenceVisitor::Run 消费它，conditional 携带按分支的绑定（dynamic_parameter_binding for conditional branch）。

第 2 层 — 图内 map：`DynamicDimensionInference`

该分析维护 map<DynamicDimension{HloInstruction*, ShapeIndex, dim} -> HloInstruction*> — 保存每个维度运行时大小的 S32 SSA 值。它由两个有序树容器支撑（一个以 HloInstruction* 为键，另一个通过 operator< @ 0x1e3a8520 以 DynamicDimension 结构为键）。公共 API（全部字节锚定）：

方法	VA
`Run(module, op_dynamism_support_fn, custom_call_handler, ShapeCheckMode, assertion_generator, exec_threads)`	`0x1e39ad20`
`SetDynamicSize(inst, ShapeIndex, dim, size_inst)`	`0x1e38f1e0`
`GetDynamicSize(inst, ShapeIndex, dim) -> HloInstruction*`	(const variant) `0x1e39bbe0`
`GetDynamicSizes(inst, ShapeIndex)`	`0x1e39bc00`
`GetDynamicShape(inst)`	`0x1e39b980`
`ForwardDynamicSize(inst, new_inst, ShapeIndex)`	`0x1e39b580`
`ReplaceAllDynamicDimensionUsesWith(a, b)`	`0x1e3927a0`
`CopyMapping(from, to, replacement_map)`	`0x1e3982c0`
`AnalyzeDynamicDimensions()`	`0x1e39b0a0`

NOTE — Run 签名的反编译交叉检查。 0x1e39ad20 处的 mangled 符号 demangle 为 Run(HloModule*, std::function<OpDynamismSupport(HloInstruction*)>, std::function<bool(HloInstruction*, DynamicDimensionInference*)>, ShapeCheckMode, std::function<void(HloInstruction*)> const&, absl::flat_hash_set<string_view> const&)。第一个 callback 是按 op 的动态性支持查询（OpDynamismSupport）；第二个是 custom_call_handler（注册的按 custom-call 的动态性推断器，接收指令和推断对象）；第四个 callback 是 assertion_generator；最后的 flat_hash_set<string_view> 是执行线程集合。这与来自 DynamicPadder::RunImpl 的调用（约第 1236 行）完全匹配。[Confidence: CONFIRMED.]

Run 从入口布局构建 DynamicParameterBinding，运行 DynamicDimensionInferenceVisitor 在每条指令上向前传播维度大小 SSA 值，然后调用 AnalyzeDynamicDimensions。

第 3 层 — 传播：`DynamicDimensionInferenceVisitor`（37 个 visitor override + 3 个 helper）

前向 visitor（Run @ 0x1e3984c0）根据操作数维度大小计算每个输出维度大小 SSA。每个 handler 接收签名为 (HloInstruction* inst, ShapeIndex, long dynamic_dim, long operand_dim, HloInstruction* dynamic_size) 的 callback。二进制导出 37 个单参数 Handle*(HloInstruction*) visitor override，外加 3 个多参数 helper（HandleDynamicConvolutionForward(HloInstruction*, long, long, HloInstruction*)、HandleDynamicConvolutionInputGrad(HloInstruction*, long, long)、HandleDynamicWindowSamePadding(HloInstruction*, HloInstruction*, long, long)），它们从 HandleCustomCall 针对三个 DynamicConvolution* custom-call 调用 — 总计 40 个 Handle* 成员函数。带字节锚点的代表性子集：

Handler	VA	传播规则
`HandleParameter`	`0x1e39a6c0`	从 `DynamicParameterBinding` 播种
`HandleSetDimensionSize`	`0x1e3929e0`	绑定维度 → operand-1（S32 大小）
`HandleGetDimensionSize`	`0x1e3923a0`	物化已存储的大小 SSA
`HandleBroadcast`	`0x1e38fae0`	通过 broadcast dims 映射；`Broadcast input and output dynamism mismatch`
`HandleConstant`	`0x1e38fba0`	清除动态性（常量是静态的）
`HandleConcatenate`	`0x1e391c20`	在 concat 维度上对操作数动态大小求和
`HandleReshape`	`0x1e394980`	分解维度组；按组穿线大小
`HandleDynamicReshape`	`0x1e3947e0`	使用显式维度大小操作数
`HandlePad`	`0x1e391360`	大小与 pad config 组合
`HandleDot` / `HandleConvolution`	`0x1e3918a0` / `0x1e391b60`	传播 batch/spatial 动态维度
`HandleDynamicConvolutionForward`	`0x1e392d00`	`GetWindowedOutputSize` 边界
`HandleDynamicWindowSamePadding`	`0x1e393060`	same-padding 边界
`HandleWhile`	`0x1e398900`	通过循环携带的 tuple 穿线动态大小
`HandleConditional`	`0x1e395d00`	按分支绑定
`HandleCustomCall`	`0x1e390580`	除非注册了 handler，否则 `Dynamic inferencing on custom call %s is not supported`
`PassThroughDynamicDimension`	`0x1e3935e0`	通用前向 helper（elementwise、select、clamp、transpose、slice…）

visitor 还通过 RequiresPadToStatic(inst, ShapeIndex) (0x1e39a7e0) 和 InsertPadToStaticOnInstruction(inst) (0x1e390920) 决定静态边界放在哪里：当下游 op 无法消费动态操作数时，插入 PadToStatic custom-call。ShapeCheckMode enum 选择是在编译期验证动态维度一致性，还是延迟到由 assertion_generator callback 发出的运行时断言；运行时模式锚点是 dynamic dimensions size %d did not match number of dimensions %d。

DynamicPadder Pass

xla::DynamicPadder::RunImpl(HloModule*, exec_threads) @ 0x16998ca0 是主动态形状 lowering pass — 开源 XLA DynamicPadder，由 DynamicPadderOptions 配置，并通过 AddPass<DynamicPadder>(DynamicPadderOptions) 添加。它位于 PreOptimization 第 21 步，紧跟在 DynamicDimensionSimplifier（第 20 步）和 ConditionalCanonicalizer 之后，而 DynamicIndexSplitter 更早位于第 4 步（见 hlo-pre-passes.md，其中记录为 2× AddPass）。

从反编译调用序列恢复的 RunImpl 流程：

构建动态性信息。 调用 DynamicDimensionInference::Run(module, op_dynamism_support_fn, custom_call_handler, shape_check_mode, assertion_generator, exec_threads)（约第 1236 行）。op_dynamism_support_fn 为每个 op 返回一个 OpDynamismSupport；锚点 op_support != OpDynamismSupport::kNoSupport 决定某个 op 保留其动态 shape，还是必须被填充。TPU tpu_compile_op_support.cc 提供按 op 的支持表。
重写无法消费动态操作数的 op。 所有重写器都位于 dynamic_padder.cc 匿名命名空间：
- PadWithScalar(inst, dim, dynamic_size, pad_value) @ 0x169a2fe0 — 用单位元标量（add-reduce 用 0，max-reduce 用 −inf，…）填充超出运行时大小的有界区域，使填充 lane 不扰动结果。
- GenerateBinaryMask(inst, dim, dims, dyn_sizes, iota, lt, is_lower) @ 0x169a6360 — 构建 iota < dynamic_size 比较 mask（该 pass 输出不变量的“padding mask”）。
- RewriteDynamicReshape(inst, ddi) @ 0x169a11c0（+ RewriteDynamicReshapeSingleGroup @ 0x169a3600）— 将 reshape 拆分为维度组并穿线动态大小。
- RewriteInputWithDynamicPadding(inst, operand, pad, dims, Window*, size_fn) @ 0x169a6ec0 — 填充 conv/window 输入，使用 GetWindowedOutputSize (0x1e3a93e0) / GetWindowedInputGradSize (0x1e3a9b40) 计算边界。由三个 DynamicConvolution{Forward, InputGrad, KernelGrad} custom-call 分支驱动。
- RewriteDynamicBinaryOp / RewriteDynamicSort（lambda），以及 reduce-window/select-and-scatter same-padding 重写器。
重写器通过 MakePadHlo (0x1e3e5560)、CreateSlice (0x1e593160)、CreateDynamicSlice (0x1e5947e0)、CreateDynamicUpdateSlice (0x1e594860)、CreateReshape (0x1e594de0) 等构建 HLO。
插入两个边界 custom-call。
- PadToStatic — 在动态值必须变成静态形状值的每一点（进入仅静态区域）插入。由 InsertPadToStaticOnInstruction / RequiresPadToStatic 插入。
- SliceToDynamic — 反向操作，在静态值必须重新标注为动态的每一点（输出边界）插入。
目标：jellyfish::dynamic_padding_handler::kPadToStatic / kSliceToDynamic，通过来自 custom_ops/dynamic_padding_handler.cc 的 CustomCallRegistration::RegisterLoweringEmitter("SliceToDynamic", dynamic_padding_emit_helper) 注册。
对确实原生支持动态的 op 去填充。 (anon)::DynamicShapeRemovingVisitor::ConvertToDynamic(inst) @ 0x169a7bc0（+ ConvertOperandsToDynamic @ 0x169a8ac0）遍历 OpDynamismSupport 为“supported”的 op，并剥离 DynamicPadder 原本会插入的 PadToStatic，避免不必要的 pad/slice 往返。镜像锚点是 Input to RemoveDynamicShapeMetadataIfPresent should be static。

NOTE — 反编译交叉检查。 已确认 DynamicPadder::RunImpl（5,918 行反编译代码）按顺序调用 Shape::is_static（约第 1071 行的 already-static early-out）、DynamicDimensionInference::Run（约 1236）、RewriteDynamicReshape（约 1410），并在各处调用 GetDynamicSize/ForwardDynamicSize/GetDynamicSizes/HasDynamicDimension；在三个 DynamicConvolution* custom-call 分支内调用 PadWithScalar 和 RewriteInputWithDynamicPadding（约 3098/3172/3332/3381/3528），并在尾部调用 DynamicShapeRemovingVisitor::ConvertToDynamic（约 5178）。ret_check kernel->shape().is_static()（约 3267）守卫 conv-kernel 路径。[Confidence: CONFIRMED.]

配套项

xla::DynamicIndexSplitter::RunImpl @ 0x164ae740（第 4 步）— 将 DynamicSlice/DynamicUpdateSlice 上的多维动态索引拆分为按维度的标量索引操作数，使每个索引都是单个 S32 SSA 值。
xla::DynamicDimensionSimplifier::RunImpl @ 0x164d0020（第 20 步）— 在 DynamicPadder 运行前折叠冗余的 Get/SetDimensionSize 链和 <= K 动态维度 op。

填充到静态 Tile 策略与运行时表示

DynamicPadder 之后，每个缓冲区都按其静态上界定大小，运行时大小位于预置到每个动态缓冲区前面的 1024 字节元数据前缀中。

缓冲区定大小

Target::ShapeWithMetadataSizeBytes(shape) (0x1d619f20) 委托给 TransferSizeUtil::ShapeWithMetadataSizeBytes (0x1d6aea00)。反编译函数体：

// xla::jellyfish::TransferSizeUtil::ShapeWithMetadataSizeBytes  @0x1d6aea00
if (element_type == 13)            // TOKEN/opaque: no payload
    prefix = 0;
else if (Shape::is_static(shape))  // fully static: no prefix
    prefix = 0;
else {                             // dynamic: prefix from layout, default 1024
    prefix = shape.layout().dynamic_shape_metadata_prefix_bytes();
    if (prefix == 0) prefix = 1024;
}
return prefix + ShapeSizeBytesRaw(shape);   // ShapeSizeBytesRaw uses static extents
```text

因此物理字节大小是 `ShapeSizeBytesRaw(shape) + prefix`，其中 `ShapeSizeBytesRaw` 根据**静态（上界）extent** 计算 — 绝不是运行时值。前缀默认是 1024，并且可以通过 `xla::Layout::dynamic_shape_metadata_prefix_bytes()` 字段按缓冲区覆盖。

> **NOTE — 前缀可由 layout 覆盖，不是扁平常量。** 动态分支读取 `layout().dynamic_shape_metadata_prefix_bytes()`，只有当该字段为零时才回退到 `1024`；`TOKEN` 元素类型 (13) 既没有前缀也没有 payload。常量 1024 来自下面的 `Target::DynamicShapeMetadataPrefixBytes()`，它负责填充 layout 字段。[Confidence: CONFIRMED — 读取自 `0x1d6aea00`。]

### 前缀常量

`Target::DynamicShapeMetadataPrefixBytes()` @ `0x1d61c4e0` 完整内容为：

```c
__int64 xla::jellyfish::Target::DynamicShapeMetadataPrefixBytes(Target *this) {
  return 1024;
}

1024 字节，适用于每一代 TPU。断言的不变量（逐字 CHECK 字符串）：

b.target().DynamicShapeMetadataPrefixBytes() % sizeof(int32_t) == 0 — 1024 = 256 个四字节 int32 维度大小槽。
b.target().DynamicShapeMetadataPrefixBytes() <= b.target().ChunkSizeBytes() — 前缀适配在一个 HBM/VMEM chunk 内。
metadata_offset == target.DynamicShapeMetadataPrefixBytes() — 数据区域正好从前缀之后开始。

Target::DynamicShapeMetadataPrefixShape() @ 0x1d61c500 构建一个维度大小的 1-D S32 shape — 前缀的逻辑类型。

设备上缓冲区布局

text

+------------------------------+--------------------------------------------+
| metadata prefix (1024 bytes) | data, padded to the static upper bound     |
| = DynamicShapeMetadataPrefix |  (= ShapeSizeBytesRaw(shape) bytes)         |
|   Bytes() = 0x400            |                                            |
| up to 256 int32 dim sizes    |  physical extent independent of runtime    |
+------------------------------+--------------------------------------------+
```text

`SliceToDynamic` 写入前缀；`PadToStatic` 读取前缀。

### Layout / tiling / MSA 处理

由于 `DynamicPadder` 在布局之前运行，后续所有 pass 都在静态数组 shape 上操作（按维度的 `is_dynamic` 位保留用于 bookkeeping）：

- **Layout** ([layout-assignment.md](layout-assignment.md))：`TpuLayoutAssignment` 通过 `ChooseCompactLayoutForShape` 从静态（上界）维度选择 layout。`dynamic_shape_metadata_prefix_bytes()` 字段随 `xla::Layout` 携带（锚点：`input_shape.layout().dynamic_shape_metadata_prefix_bytes() is expected to be non-zero, where input_shape =`），记录该缓冲区携带前缀。
- **Tiling** ([loop-tiling-unrolling.md](loop-tiling-unrolling.md))：tiling 使用静态 extent。`lowering_util::DynamicShapeSizeCompactRaw` (`0x1c6ca220`) 和 `DynamicShapeSizeCompactForDmaRaw` (`0x1c6ca8a0`) 在 DMA 时从维度大小标量计算*实际*运行时传输字节数，因此 DMA 只搬运 live 区域，而不是完整的填充到上界缓冲区。
- **MSA**（memory-space assignment）：为每个缓冲区分配 `Target::ShapeWithMetadataSizeBytes(shape)` — 填充到上界 + 前缀。断言 `allocation->size() == target.ShapeWithMetadataSizeBytes(allocation->shape())`。MSA 从不看到运行时定大小的缓冲区。

TF/host 边界使用 `tensorflow::XlaTpuPaddedShapeFn(TpuTopology, Shape, Shape*)` @ `0xf7d1cc0`（→ `TransferSizeUtil::SetPaddedShape`）为动态 XLA shape 计算设备上的 padded shape（将每个动态维度填充到其上界，并遍历 tuple）。

---

## DynamicSlice / DynamicUpdateSlice / Reshape Lowering

`DynamicPadder` 之后，模块是静态的，并带有两个边界 custom-call。`dynamic-pad` 和 `dynamic-reshape` 永远不会存活到 LLO — `DynamicPadder` 在 HLO 阶段将它们完全重写为 `Pad` + `DynamicSlice`/`Reshape` + mask 序列。剩余的一等动态索引 op 由 TPU HLO→LLO `LoweringEmitter` 直接 lowering：

| Emitter | VA | 作用 |
|---|---|---|
| `jellyfish::LoweringEmitter::HandleDynamicSlice(hlo)` | `0x10c3b0c0` | 带运行时基址偏移的静态 extent 索引 load |
| `jellyfish::LoweringEmitter::HandleDynamicUpdateSlice(hlo)` | `0x10c3b640` | 静态 extent 索引 store |
| `jellyfish::DynamicUpdateSliceEmitter` | `0x10c66ec0` | DUS 变体（`OpEmitter::Emit<>`） |
| `jellyfish::AsyncDynamicIndexEmitter` | `0x10c415e0` | pipeline 形式 — 在一个 bundle 中计算索引标量，将索引 load/store 延迟到后续 bundle |

由于 slice extent 是静态的（slice 进入一个填充到上界的缓冲区），lowering 是带运行时基址偏移的静态 extent 索引 copy，基址偏移由 `DynamicIndexSplitter` 产生的 S32 索引标量计算。X128/precision 和 sharding visitor 复用相同 handler：`XPrecisionRewriteVisitor::HandleDynamicSlice/UpdateSlice` (`0x1115fb40`/`0x1115fd00`) 和 `DimLabelPropagation::HandleDynamicSlice/UpdateSlice` (`0x11197ca0`/`0x11198220`)。peephole `SliceToDynamicCopyMover` (`0x10fc0240`) 将 `Copy` 推过 `SliceToDynamic(Copy(x))`，以消除动态边界处的冗余 copy。

### 两个边界 emitter

它们是运行时契约的核心。

**`PadToStaticEmitter::Emit()` @ `0x10c9ad40`** 接收一个动态形状输入缓冲区（前缀 + 填充到上界的数据），并生成 {静态形状数据数组, S32 维度大小标量} 的 tuple。已确认的 LLO 序列：

```text
DynamicShapeMetadataPrefixBytes()                 // locate the 1024-byte prefix
  -> CHECK prefix % sizeof(int32_t) == 0
LloRegionBuilder::Vld(...)                          // vector-load dim-size metadata
lowering_util::ScalarToSreg(...)                    // move size into a scalar reg
lowering_util::ComputeBoundsInChunks(sizes, b)      // dim sizes -> chunk counts
lowering_util::CalculateDynamicCompact2ndMinorRatio // re-derive compact tiling ratio
Target::ChunkCountsWithTmp(shape, tmp)              // pad-to-bound chunk counts
PipelineEmitter::SetDynamicIterationBounds(...)     // iteration count driven by runtime size
PipelineEmitter::Emit(...)                          // software-pipelined chunk transfer

NOTE — 反编译交叉检查。 上述所有内容均直接从 0x10c9ad40 读取：DynamicShapeMetadataPrefixBytes() & 3 可整除检查及其匹配的 CHECK 字符串（约第 274/277 行）、Vld（约 291）、ScalarToSreg（约 446）、ComputeBoundsInChunks（约 499）、ChunkCountsWithTmp（约 270/271）、CalculateDynamicCompact2ndMinorRatio（约 705/841），以及 PipelineEmitter::SetDynamicIterationBounds（约 630）。[Confidence: CONFIRMED.]

SliceToDynamicEmitter::Emit() @ 0x10c9c6c0 是反向操作：它接收静态形状数据数组 + S32 维度大小标量，写入前缀，并生成动态形状缓冲区。已确认的 LLO 序列：

text

DynamicShapeMetadataPrefixShape()                   // the 1-D S32 prefix shape
  + DynamicShapeMetadataPrefixBytes() (% 4 check)
LloRegionBuilder::Vlaneseq / VimmS32 / VeqS32 / Vselect   // build per-lane validity mask
lowering_util::BroadcastScalarToVreg(...)            // broadcast each dim size
deep_copy_util::MemsetInGranules(...)                // zero/init the prefix region
lowering_util::ComputeBoundsInChunks(...)            // sizes -> chunk counts
lowering_util::CalculateDynamicCompact2ndMinorRatio  // compact tiling ratio
... CopyArray after the prefix
```text

> **NOTE — 反编译交叉检查。** 读取自 `0x10c9c6c0`：`Vlaneseq`/`VimmS32`/`VeqS32`/`Vselect` mask 构造（约第 291–317 行）、`BroadcastScalarToVreg`（约 314）、前缀 `% sizeof(int32_t)` 检查（约 324/327）、`DynamicShapeMetadataPrefixShape`（约 336）、`ComputeBoundsInChunks`（约 430）、`MemsetInGranules`（约 541）和 `CalculateDynamicCompact2ndMinorRatio`（约 613/744）。[Confidence: CONFIRMED.]

---

## 不支持和被拒绝的情况

TPU 明确拒绝*无界*动态性以及若干 op 特定的动态情况。逐字字符串锚点：

| 锚点 | 含义 |
|---|---|
| `unbounded dynamism is not supported` / `Unbounded dynamism is disabled for instruction: %s` | 只有有界动态维度可编译 |
| `AllToAll does not support bounded dynamic shapes` / `AllToAllTuple does not support unbounded dynamic shapes` | collective 拒绝动态性 |
| `CustomCall "%s" is not supported to have a dynamic dimension` | 大多数 custom-call 必须是静态形状 |
| `Dynamic inferencing on custom call %s is not supported` | 没有注册的动态性 handler |
| `bitcast-convert is not valid for dynamic shape %s->%s` | bitcast-convert 拒绝动态性 |
| `The output of iota must not have dynamic dimensions` | iota 输出必须是静态的 |
| `Dynamic shapes are not supported for host buffers` / `dynamic shapes not supported in allocations` | host/pinned allocation 拒绝动态性 |
| `MemRefType don't support dynamic shapes` | MLIR memref 必须是静态的 |

StableHLO 无界 `Dynamic*Op` 家族（`DynamicReshapeOp`、`DynamicBroadcastInDimOp`、`DynamicConvOp`、`DynamicGatherOp`，…）在二进制中存在 verifier，但尚未确认它们是否在 TPU 上可达，还是仅作为前端 import artifact，随后被 `unbounded dynamism is not supported` 拒绝。[Confidence: LOW.]

---

## 未恢复的内容

- **1024 字节前缀的精确内部字节/维度顺序。** 推断为一个维度大小的 1-D S32 数组（`DynamicShapeMetadataPrefixShape`）；尚未从 live buffer 字节转储确认顺序是 major-to-minor 还是 minor-to-minor，以及 tuple 子 shape 是共享前缀还是各自拥有前缀。[Confidence: MEDIUM.]
- **完整 `DynamicPadderOptions` proto 字段集。** 通过 `Run` 签名 + 锚点确认的字段有：`shape_check_mode`、`op_support_from_compute`（`OpDynamismSupport` fn）、`assertion_generator`、`slice_dynamic_output`。完整 `dynamic_padding.proto` descriptor 尚未逐字段解码。[Confidence: MEDIUM.]
- **按 op 的 `OpDynamismSupport` 表**（`tpu_compile_op_support.cc`）— 哪些 op 保留动态操作数，哪些要求 `PadToStatic` 边界，尚未枚举。[Confidence: LOW.]
- **每个 `DynamicDimensionInferenceVisitor::Handle*` 的按 handler 算术。** 开源传播算法已知，TPU 二进制遵循它；handler 函数体做过抽样，未穷尽追踪。[Confidence: MEDIUM.]
- **SparseCore 动态维度子系统**（`LowerDynamicDimensionSizePass`、`sc_tpu.set/get_dynamic_dimension_size`、`DynamicBoundedSlicedInput`、`ConvertStaticToDynamicEmitter`）— 这是一个独立的 MLIR 级路径，不在本页范围内。
- **Conv/window same-padding 边界算术**（`GetWindowedOutputSize` / `RewriteDynamicConvolution*`）— 本页锚定，但未追踪到重新实现精度。

---

## 交叉引用

- [hlo-pre-passes.md](hlo-pre-passes.md) — 有序预 pass 表，包含 `DynamicIndexSplitter` (#4)、`DynamicDimensionSimplifier` (#20) 和 `DynamicPadder` (#21, 2× `AddPass`)；本页是这些行的算法细节。
- [compile-phases.md](compile-phases.md) — 顶层阶段顺序；`DynamicPadder` 在 PreOptimization 阶段内、layout/MLIR 之前运行。
- [overview.md](overview.md) — 编译器导览以及 HLO → … → LLO IR 层级栈；负责边界 emitter 馈入的 MLIR handoff。
- [layout-assignment.md](layout-assignment.md) — `TpuLayoutAssignment` 携带 `dynamic_shape_metadata_prefix_bytes()` layout 字段；它在 `DynamicPadder` 之后看到的是静态模块。
- [loop-tiling-unrolling.md](loop-tiling-unrolling.md) — tiling 使用静态 extent；`DynamicShapeSizeCompactForDmaRaw` 在运行时将 DMA 定大小到 live 区域。
- [back to index](../index.md)

动态形状支持 ​

摘要 ​

有界动态 Shape 模型 ​

第 2 层 — 图内 map：DynamicDimensionInference ​

第 3 层 — 传播：DynamicDimensionInferenceVisitor（37 个 visitor override + 3 个 helper） ​

DynamicPadder Pass ​

配套项 ​

填充到静态 Tile 策略与运行时表示 ​

缓冲区定大小 ​

设备上缓冲区布局 ​

动态形状支持

摘要

有界动态 Shape 模型

第 2 层 — 图内 map：`DynamicDimensionInference`

第 3 层 — 传播：`DynamicDimensionInferenceVisitor`（37 个 visitor override + 3 个 helper）

DynamicPadder Pass

配套项

填充到静态 Tile 策略与运行时表示

缓冲区定大小

设备上缓冲区布局