SFLAG 同步标志内存协议

本页所有地址均适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（构建 libtpu_lts_20260413_b_RC00，build-id md5 89edbbe81c5b328a958fe628a9f2207d）。该映像未剥离符号；反混淆后的 C++ 符号名按原文逐字引用，且 .text VMA == 文件偏移。其他版本会有所不同。

摘要

SFLAG 是 TPU 的原子同步层级：一个小型、按字粒度的寄存器文件（MemorySpace::kSflag = 6，区别于 HBM/VMEM/CMEM/SMEM — 见内存层次概览），每个引擎都会在其上自旋，用来协调 DMA、跨引擎交接、collective 和 barrier。flag 通过每核心表中的扁平同步标志编号寻址；编译器从不直接触碰 SFLAG 字节，而是用 LloRegionBuilder::SflagImmPtr(number, name) 构造一个 LloValue 指针，然后通过固定的一族 Vsync* / Vwait* region-builder 原语驱动它，这些原语降低为 mlir::llo 的 VSync*Op / VWaitGeOp 指令类。

本页只负责四件事。(1) 内存模型： SFLAG 是 32-bit 字的扁平数组；SflagImmPtr 把 flag 编号 n 转成一个带 kSflag 标签的指针，其 byte_offset = 4·n，形状为 1 元素 S32，因此 flag 编号算术是精确且按字粒度的。(2) flag 值语义： 一个 flag 携带饱和计数器或一次性 done-bit，由每次写入时的 0x100 done bit 选择，并由 Target::HasExtraDoneBitInSyncFlags() 按代门控；硬件是否把该字解释为计数器或布尔值，由 SyncFlagCountModeBitOffset 决定（Viperfish-non-lite/Ghostlite 上是 bit 272，其他位置是 0）。(3) 驱动原语： VsyncSet（写入）、VsyncAdd / VsyncAddDone（原子递增，不带/带 done）、VwaitGeSV（阻塞直到 ≥ 阈值）及其 Eq/Ne/Lt/Done 同族、VsyncRead（把计数器 pop 到标量中），以及 remote 和 public-access 变体。(4) SflagImmPtr 本身以及它消费的按代 dummy/count 常量。

不在范围内并由其他页面负责的内容：barrier 绑定到哪个 flag 编号见 Barrier → SFLAG Number Binding；驱动这些原语的跨核心 barrier datapath 是 barrier 子系统，尤其是 tree-barrier Vsync emitter；remote-DMA sflag 地址编码器见 remote-sflag encoders。本页是这些页面站立其上的基底。


内存空间	`xla::jellyfish::MemorySpace::kSflag` = 6（operand-space 标签，由 `SflagImmPtr` 烘进指针）
字	32-bit（S32）；`byte_offset = 4 · sync_flag_number`
指针 ctor	`LloRegionBuilder::SflagImmPtr(int number, string_view name)` @ `0x1d5185a0`
BarnaCore 指针 ctor	`LloRegionBuilder::BarnaCoreSflagImmPtr(int, string_view)` @ `0x1d538400`
Set	`LloRegionBuilder::VsyncSet(LloValue* sf, LloValue* v, optional<bool> done)` @ `0x1d54cce0` → `CreateVectorSyncFlagSet`
Add	`LloRegionBuilder::VsyncAdd(sf, v)` @ `0x1d523200` → `CreateVectorSyncFlagAdd(flag=0)`
Add+done	`LloRegionBuilder::VsyncAddDone(sf, v)` @ `0x1d54e380` → `CreateVectorSyncFlagAdd(flag=257)`
Wait ≥	`LloRegionBuilder::VwaitGeSV(sf, thr, bool yieldable)` @ `0x1d522f80` → `CreateVectorWait(391 + 8·yieldable)`
Read	`LloRegionBuilder::VsyncRead(sf)` @ `0x1d524220` → `PushSyncFlagToFifo` + (`PopSfrFifo` \| `CreateScalarV2SPop`)
Dummy 编号	`Target::GetDummySyncFlagNumber()`（vtable `+0x4F8`/idx 159）— JF = 7，PF/VF/GL = 0
Count-mode bit	`Target::SyncFlagCountModeBitOffset()` — 272（VF-non-lite/GL），0（JF/PF/VF-lite）
Done-bit 支持	`Target::HasExtraDoneBitInSyncFlags()`（按代谓词）
置信度	CONFIRMED（按字节锚定），除非某行或 callout 另有说明

1. 内存模型 — 32-bit 字的扁平数组

目的

重新实现者寻址 SFLAG 只需要一个事实：同步标志编号是 32-bit 字数组中的索引，flag n 的字节偏移是 4·n。没有按 flag 的 descriptor，没有对齐空隙，指针层级没有 sub-word packing；确实存在的 packing（§4）位于指令的操作数编码中，而不是地址中。编号空间本身（哪些编号保留、哪些可分配）由 Target 对象上的分区固定，并且 barrier 槽位记录在 Barrier → SFLAG Number Binding；这里我们只需要知道一个编号会解析为一个字地址。

`SflagImmPtr` — 编号 → 指针

LloRegionBuilder::SflagImmPtr @ 0x1d5185a0 是 SFLAG LloValue 指针的唯一构造器。反编译后，它构造一个经过验证的 1 元素 S32 shape，并以字节偏移和 SFLAG operand-space id 调用通用 ImmPtr：

function SflagImmPtr(builder, int number, string_view name):    // 0x1d5185a0
    shape = ShapeUtil::MakeValidatedShape(/*element_type=*/4,    //   4 == S32 (xla::PrimitiveType)
                                          /*dims=*/0, /*minor_to_major=*/0)
    //                  byte_offset       shape  space  name
    return builder.ImmPtr(4 * number,     shape, 6,     name)    //   space 6 == MemorySpace::kSflag
```text

两个数字很核心，并且都逐字节确认：

- **Element type `4` = S32。** 每个同步标志都是 32-bit 整数字。shape 是 rank-0（一个标量），因此 SFLAG 指针恰好命名一个字。
- **`byte_offset = 4 · number`**（`4LL * a2` 参数）。这个四倍因子*就是*字步幅；这解释了为什么 `SflagWordSizeBytes()`（`Target+0x504`，见[概览的分配器矩阵](overview.md#3-the-per-space-allocator--alignment-matrix)）是 SFLAG 的粒度和对齐，也解释了为什么缓存 `SflagWordSizeLog2()`（`Target+0x4c8`）以便地址算术能用移位而不是乘法。

空间参数 `6` **就是** `MemorySpace::kSflag` 本身；`SflagImmPtr` 把 operand-space 标签直接盖印进指针。这与 §3 中每个原语检查的同一个 `6` 一致（`(*((_BYTE*)sf+11)>>2)&0x1F == 6`），并且与结构相同的 `BarnaCoreSflagImmPtr` @ `0x1d538400` 逐字节确认，后者传入空间 `0xA`（= `kBarnaCoreSflag` = 10）。这里没有单独的“render id”被烘入：descriptor 路径为 sflag 渲染的 *DMA driver-resource* id 是另一个完全不同的数字（`sflag(6) → 0`），只在 descriptor 边界由 `MemorySpaceToDriverResource` 产生，绝不由此构造器产生（见[枚举校正](overview.md#2-the-memoryspace-enum)）。

> **注意 —** `BarnaCoreSflagImmPtr` @ `0x1d538400` 是 **BarnaCore** 同步标志层级的结构相同构造器（`MemorySpace::kBarnaCoreSflag` = 10；构造器向 `ImmPtr` 传入空间 `0xA`）。它是物理上独立的寄存器文件，拥有自己的大小 accessor（`Target::BarnaCoreSflagSizeBytes()`，`Target+0x478`，由 `HasBarnaCore` vtable 谓词保护）。面向 BarnaCore flag 的代码必须使用这个构造器；这两个指针家族绝不可互换。

### 为什么是指针，而不是编号

§3 中每个驱动原语都接收一个 `LloValue*` 同步标志操作数，而不是裸整数。原因是 memory-space CHECK：每个原语断言 `sync_flag->memory_space() == MemorySpace::kSflag`（下文解码），所以 flag 必须以已类型化的 SFLAG 指针到达。该检查从 `LloValue` header 读取操作数的空间标签（`(*((_BYTE*)sf + 11) >> 2) & 0x1F`），并通过 `LloCheckForFailure<MemorySpace, MemorySpace, LloCheckOp::Eq>` 与常量 `kSflag` 比较。由 `SflagImmPtr` 构造的指针会通过；任何其他操作数都会使编译失败，并给出 `"sync_flag->memory_space() == MemorySpace::kSflag"` 以及出错值的 mnemonic dump。

---

## 2. Flag 值语义 — counter vs. done-bit

### 目的

对系统其余部分而言，SFLAG 字不是普通整数：它被解释为**饱和计数器**（生产者 add，消费者等待阈值）或**一次性 done-bit**（生产者标记 done，消费者 wait-done）。采用哪种解释会被*选择两次*：一次在每次写入时选择（指令级 done bit），一次按代选择（硬件 count-mode bit）。弄错这一点会静默死锁：消费者对一个只设置了 done-bit 的生产者等待计数器阈值，将永远自旋。

### 两种值模式

| 模式 | 生产者 | 消费者 | 用途 |
|---|---|---|---|
| **Counter** | `VsyncAdd`（原子 +）/ `VsyncSet` 到 N | `VwaitGeSV(thr)` — 阻塞直到值 ≥ thr | N 路 fan-in：N 个生产者各自 `VsyncAdd 1`，一个消费者 `VwaitGe N` |
| **Done** | `VsyncSet(v, done=true)` / `VsyncAddDone` | `VwaitDone` | 1:1 生产者/消费者完成握手 |

### 指令级 done bit

done-bit 作为指令 flag word 的 bit `0x100`（= 256）携带。反编译精确钉住了两个生产者：

- `VsyncAdd` @ `0x1d523200` 发出 `CreateVectorSyncFlagAdd(sf, v, /*flag=*/0, …)` — flag word `0`，纯计数器递增。
- `VsyncAddDone` @ `0x1d54e380` 发出 `CreateVectorSyncFlagAdd(sf, v, /*flag=*/257, …)` — flag word `0x101` = `0x100`（done）| `0x1`（共同设置的低位）。单独的 bit `0x100` 是 done 选择器；同一个 `0x101` 掩码也作为 done 谓词出现在 `VsyncSet` 中。
- `VsyncSet` @ `0x1d54cce0` 接收一个 `optional<bool> update_done_to`。当 optional 存在时（其 `0x100` 判别位被设置），builder 在发出指令前断言目标支持它：

```c
// VsyncSet @ 0x1d54cce0 — done-bit guard (lines ~47-71)
if (update_done_to.has_value() && !target().HasExtraDoneBitInSyncFlags()):
    CHECK_FAIL("!update_done_to.has_value() || target().HasExtraDoneBitInSyncFlags()")
              << "Done bit not supported for this target."        // llo_region_builder.cc:8061
emit CreateVectorSyncFlagSet(sf, value, flag_word, …)              // line 73

陷阱 — done-bit 是按目标的能力，不是通用能力。Target::HasExtraDoneBitInSyncFlags() 是虚谓词（通过模块的 target 指针以 target().… 到达）。在返回 false 的一代上，任何带 done 值的 VsyncSet 或任何 VwaitDone 都是硬编译失败（"Done bit not supported for this target."，llo_region_builder.cc:8061）。总是发出 done-style completion 的重新实现无法在没有额外 done bit 的目标上编译；counter-mode（VsyncAdd + VwaitGe）是可移植 lowering。

硬件 count-mode bit（272）

硬件寄存器把 flag 字当作计数器还是 done-bit，由 Target::SyncFlagCountModeBitOffset() 选择；这是多字 SFLAG 寄存器编码内的一个位位置。按代解码如下：

代	`SyncFlagCountModeBitOffset()`	地址	解释
Jellyfish（v2）	0	`0x1d491380`	未使用 count-mode bit；模式隐式
Pufferfish（v4）	0	`0x1d495860`	模式隐式
Viperfish（v5）	272 non-lite / 0 lite	`0x1d49bca0`	`codename=="lite" ? 0 : 272`
Ghostlite（v6e）	272（`0x110`）	`0x1d4988c0`	count-mode bit 位于寄存器 bit 272

Viperfish accessor（0x1d49bca0）解码为：除非 codename 字符串是 4 字符 "lite"，否则返回 272；它先用 4 测试字符串长度（v2 != 4 → return 272），然后对 4 字符串把前导 dword 与 1702127980（= 0x6574696c，"lite" little-endian）比较，并且只在匹配时返回 0。因此 viperfish-lite 禁用 count-mode bit，而 viperfish-std 保留它；这是 VMEM 和 barrier 页面也观察到的同一个 lite-codename 字符串比较分叉。（这个 4 是 codename 字符串长度，不是 TpuVersion ordinal。）值 272 是位索引，不是字节偏移：它在支持每个 flag 的更宽硬件寄存器中选择“count mode vs. done mode”。

[LOW] bit 272 所在的多字 SFLAG 寄存器的完整 bit-field 映射（counter 宽度、done bit、public-access bit）未在此解码。bit 272 作为按代 accessor 返回的 count-mode 选择器已 CONFIRMED；周围字段未枚举，需要解码 VSync*Op 寄存器编码器的 LLO→ISA lowering。

3. 驱动原语 — `Vsync` / `Vwait`

目的

所有 SFLAG 变更和观察都流经一族 LloRegionBuilder 方法。它们共享固定形状：接收一个 LloValue* 同步标志指针（写入/等待还接收一个 LloValue* 值或一个 int 立即数），断言它是 kSflag，然后构造并追加匹配的 mlir::llo 指令。重新实现者通过复现这张表来复现协议；动词集合是封闭的。

原语表

Builder 方法	地址	LLO 指令	语义
`VsyncSet(sf, v, opt<bool> done)`	`0x1d54cce0`	`CreateVectorSyncFlagSet`	写入 `v`（可选设置 done）
`VsyncAdd(sf, v)`	`0x1d523200`	`CreateVectorSyncFlagAdd(flag=0)`	原子 `*sf += v`
`VsyncAddDone(sf, v)`	`0x1d54e380`	`CreateVectorSyncFlagAdd(flag=257)`	原子 add 并设置 done
`VsyncAddInGranules(sf, MemUnit)`	`0x1d54e160`	`CreateVectorSyncFlagAdd`	增加一个 granule-count（byte→word）
`VsyncRead(sf)`	`0x1d524220`	`PushSyncFlagToFifo` → pop	将 counter 读入标量寄存器
`VwaitGeSV(sf, thr, yield)`	`0x1d522f80`	`CreateVectorWait(391 + 8·yield)`	阻塞直到 `*sf ≥ thr`
`VwaitEqSV` / `VwaitNeSV` / `VwaitLtSV`	—	`CreateVectorWait`（cond ≠ Ge）	阻塞直到 `==` / `!=` / `<`
`VwaitDone(sf, yield)`	—	`CreateVectorWaitDone`	阻塞直到 done-bit 被设置
`VsyncSetRemote(sf, CoreLoc, v)`	`0x1d54e120`	`CreateVectorSyncFlagSetRemote`	在远端核心上设置 flag
`VsyncAddRemote(sf, CoreLoc, v, b)`	`0x1d522f40`	`CreateVectorSyncFlagAddRemote`	加到远端核心 flag
`VsyncAddRemoteInGranules`	`0x1d54e4e0`	`CreateVectorSyncFlagAddRemote`	远端 granule-count add
`VsyncPublicAccessSet(sf, v)`	—	`CreateVectorSyncFlagPublicAccessSet`	set + 发布给跨引擎
`VsyncMarkAsPublic` / `MarkAsPrivate(sf)`	—	（public-access bit）	切换跨引擎可见性

`VsyncAdd` — 规范生产者

function VsyncAdd(builder, LloValue* sync_flag, LloValue* value):   // 0x1d523200
    CHECK(sync_flag->memory_space() == MemorySpace::kSflag,         //   llo_region_builder.cc:8280
          "sync-flag argument was not in the sync-flag memory space: " + mnemonic(sync_flag))
    inst = LloInstruction::CreateVectorSyncFlagAdd(sync_flag, value, /*flag=*/0, module)
    return region.AppendInstruction(inst)
```text

`VsyncAddDone` 除了 `flag = 257`（`0x101`，done bit）外逐字节相同；这确认了“add”和“add-and-mark-done”是*同一条*指令，只是 flag word 不同，而不是两条指令。

### `VwaitGeSV` — 规范消费者

`VwaitGeSV` @ `0x1d522f80` 是最丰富的原语，也是 barrier wait lowering 到的原语。它按顺序做五件事：

```c
function VwaitGeSV(builder, LloValue* sync_flag, LloValue* threshold, bool yieldable):  // 0x1d522f80
    if (yieldable):
        CHECK(target().SupportsYieldableOps(), "target().SupportsYieldableOps()")   //   :8006
    CHECK(sync_flag->memory_space() == MemorySpace::kSflag,                          //   :8008
          "sync_flag->memory_space() == MemorySpace::kSflag")
    if (autoflag tracing-vwait enabled):                                            //   AutoOr<bool> probe
        emit CreateTracingVectorWait(sync_flag)                                     //   pre-wait trace hook
    inst = LloInstruction::CreateVectorWait(/*opcode=*/391 + 8*yieldable,           //   Ge condition
                                            sync_flag, threshold, /*neg=*/…, module)
    appended = region.AppendInstruction(inst)
    dummy = SflagImmPtr(target().GetDummySyncFlagNumber(), "dummy sync flag", …)     //   vtable +1272
    if (autoflag tracing-vwait enabled):
        emit CreateTracingVectorWait(dummy)                                          //   post-wait trace hook
    return appended

重新实现者必须携带三个细节：

wait opcode 是 391 + 8·yieldable。 非 yieldable 的 Ge-wait 是 opcode 391；yieldable 形式（允许 sequencer 在自旋时 yield）是 399。只有当 Target::SupportsYieldableOps()（JF/PF/VF/GL 上存在的按代谓词）返回 true 时，yieldable 才合法，否则编译会在 llo_region_builder.cc:8006 失败。
它穿过一个 dummy 同步标志。 wait 之后，builder 在 GetDummySyncFlagNumber()（作为 +1272/idx 159 处的 vtable 调用到达，命名为 "dummy sync flag"）处制造一个 SFLAG 指针。Jellyfish 上这个编号是 7；PF/VF/GL 上是 0。dummy 槽位存在是为了让 post-wait 机制（以及 tracing）始终有一个有效 flag 可命名；它与按代配置暴露的 DummySyncFlagNumber 相同（§5）。
Tracing 是条件性的。 两次 CreateTracingVectorWait 发出都由一个 AutoOr<bool> proto flag（xla::jellyfish::AutoProto）门控；禁用时会跳过它们，因此 steady-state lowering 是单个 CreateVectorWait。Eq/Ne/Lt 同族只在编码进 opcode 的条件上不同；VwaitDone 使用 CreateVectorWaitDone，并读取 done-bit 而不是比较 counter。

`VsyncRead` — 观察 counter

VsyncRead @ 0x1d524220 把当前 flag 值读入标量寄存器。它不做比较；它 pop：

function VsyncRead(builder, LloValue* sync_flag):                  // 0x1d524220
    fifo = PushSyncFlagToFifo(builder, sync_flag)                  //   stage the flag into the read FIFO
    if (target()->vtable[+640]()):                                 //   target supports SFR FIFO pop
        return PopSfrFifo(builder, fifo)
    return AppendInstruction(CreateScalarV2SPop(fifo))             //   else vector→scalar pop
```text

读取分两阶段：flag 被 push 进 sync-flag FIFO，然后要么通过专用 SFR-FIFO 路径 pop（当目标的 `+0x280`/idx 80 vtable 谓词成立），要么通过通用的 `ScalarV2SPop` vector-to-scalar 移动。无论哪种方式，消费者最终都会在 SREG 中得到 counter 值。

### Public access — 跨引擎可见性

`VsyncPublicAccessSet` / `VsyncMarkAsPublic` / `VsyncMarkAsPrivate` 切换 flag 的 *public-access* 状态，该状态控制子引擎的 flag 是否能在其拥有引擎之外被观察到。它们降低为 `CreateVectorSyncFlagPublicAccessSet`，并在 ISA 层级降低为 `TensorCoreVectorMisc_SetPublicAccess` / `_ReadSyncPublicAccess`（每个都携带 `SyncFlagNumberField`）；在 `gxc.gfc`、`gxc.glc` 和 `vxc.vfc` 命名空间中确认了按引擎的 ISA 实现。让一个子引擎命名另一个子引擎 flag 的详细跨模式地址空间转换，是 [remote-sflag encoder](../barrier/remote-sflag-encoders.md) 的故事；这里知道 public bit 是可见性门即可。

> **注意 —** `*Remote` 变体（`VsyncSetRemote` @ `0x1d54e120`，`VsyncAddRemote` @ `0x1d522f40`）接收一个 `CoreLocationBase` 参数并面向*不同*核心上的 flag，降低为 `…Remote` 指令类。它们是 ICI collective ack 的构建块（remote DMA 写会自动递增目的端 flag）。remote flag 的地址编码由 [remote-sflag encoders](../barrier/remote-sflag-encoders.md) 负责。

---

## 4. 操作数打包

### 目的

指针层级模型（§1）是每个 flag 一个字。*操作数*编码更紧凑，因为 flag 值很小且 ISA 会打包它们。只读 §1 的重新实现者会错误估计馈入原语的标量操作数大小。

### 每个标量寄存器两个 flag

携带同步标志操作数的 SC 标量寄存器会把**两个** flag 编号打包进一个 32-bit SREG（低/高半部），由断言 `num_sync_flags_encoded_per_sreg == 2` 确认。因此，一个按编号命名 flag 的原语会消耗半个 SREG；字节层级的地址算术（`4·n`）不受影响，但操作数寄存器记账必须使用 2-per-sreg packing。

### 按字粒度跟踪

SC 分配高水位属性以**字**而不是字节记录 SFLAG 占用：`AllocationHighWaterMarkAttr::getSflagWords()` @ `0x1458fec0`。这与 `SflagImmPtr` 的 `4·n` 步幅一致；分配器计数 flag（字），寻址器乘以四。

### 值中的 done bit

如 §2 所述，生产者侧的 done 选择器是指令 flag word 的 bit `0x100`（`VsyncAddDone` flag `0x101`，`VsyncSet` 的 `update_done_to` 判别位）。这是操作数层级的 packing，不是地址层级：同一个 flag 字，同一个 `4·n` 地址，只是控制位不同。

---

## 5. 按代 SFLAG 常量

### 目的

驱动原语消费少数按代常量：dummy flag 编号、count-mode bit，以及 done-bit / yieldable 能力。这些是协议仅有的按代输入；其他一切（`4·n` 步幅、kSflag CHECK、动词集合）都与代无关。

### 常量

| 常量 / 能力 | Accessor | JF（v2） | PF（v4） | VF（v5） | GL（v6e） |
|---|---|---:|---:|---:|---:|
| Dummy flag 编号 | `GetDummySyncFlagNumber()` | **7** | 0 | 0 | 0 |
| Count-mode bit offset | `SyncFlagCountModeBitOffset()` | 0 | 0 | **272**（lite: 0） | **272** |
| flags 中的 done-bit | `HasExtraDoneBitInSyncFlags()` | 按代谓词 | 按代 | 按代 | 按代 |
| Yieldable waits | `SupportsYieldableOps()` | 按代谓词 | 按代 | 按代 | 按代 |
| Remote-DMA flag 上限 | `GetSyncFlagNumberLimitForRemoteDma()` | `jxc::Max(=59)+1` / `+0x530` | `Target+0x530` | `Target+0x530` | `Target+0x530` |

按代 override 地址（已解码立即数）：`JellyfishTarget::GetDummySyncFlagNumber` @ `0x1d4908c0`（= 7）；`Pufferfish/Viperfish/Ghostlite::GetDummySyncFlagNumber` @ `0x1d494ea0` / `0x1d49b280` / `0x1d4980c0`（全 = 0）。`SyncFlagCountModeBitOffset`：`0x1d491380`（JF = 0）、`0x1d495860`（PF = 0）、`0x1d49bca0`（VF：`0x110` non-lite / 0 lite）、`0x1d4988c0`（GL = `0x110`）。`SupportsYieldableOps` 存在于全部四代（`JellyfishTarget` … `GhostliteTarget`）。

> **注意 —** Jellyfish dummy 编号是 **7** 而每个后续代都使用 **0**，这正是会破坏朴素重新实现的“设计上的 off-by”怪异点：Jellyfish 上的 `VwaitGe` 会穿过 dummy flag 7，因此 flag 7 必须在 JF 上保留为 dummy，而在 PF/VF/GL 上是空闲的。Jellyfish 上 `GetSyncFlagNumberLimitForRemoteDma()` 的上限 59（`asic_sw::deepsea::jxc::MaxSyncFlagNumberForRemoteDma` @ `0x1d62da80` = `0x3b`）意味着 JF 上只有 flag 0..59 可作为 remote-DMA completion target；编码后的 remote-sflag 字段受限于该范围。
>
> **[LOW]** 字面按 codename 的 SFLAG 字节大小 / 字大小（`SflagSizeBytes()` @ `Target+0x468`，`SflagWordSizeBytes()` @ `Target+0x504`）由嵌入的 `chip_parts.binarypb` 在启动时填充，且**不能**从 `.text` 静态提取。accessor *偏移*和 `4·n` 步幅已 CONFIRMED；数字大小是与 [VMEM allocator](vmem-allocator.md) 和[概览](overview.md)共享的 memfile 依赖。

---

## 6. 验证记录

> 针对此页，从 `libtpu.so` v0.0.40 的 IDA 反编译逐字节重新推导：
>
> - `SflagImmPtr` @ `0x1d5185a0`：`MakeValidatedShape(element_type=4 == S32)`；`ImmPtr(byte_offset = 4·number, shape, space=6, name)` — 字步幅和 render space id 精确。
> - `VsyncAdd` @ `0x1d523200`：CHECK `sync_flag->memory_space() == MemorySpace::kSflag`（`llo_region_builder.cc:8280`）；`CreateVectorSyncFlagAdd(…, flag=0, …)` — 纯 counter add。
> - `VsyncAddDone` @ `0x1d54e380`：相同 CHECK（`:8289`）；`CreateVectorSyncFlagAdd(…, flag=257=0x101, …)` — done bit `0x100`。
> - `VsyncSet` @ `0x1d54cce0`：CHECK kSflag（`:8057`）；done-bit guard `!update_done_to.has_value() || target().HasExtraDoneBitInSyncFlags()`（`:8061`，`"Done bit not supported for this target."`）；`CreateVectorSyncFlagSet`。
> - `VwaitGeSV` @ `0x1d522f80`：yieldable guard `target().SupportsYieldableOps()`（`:8006`）；CHECK kSflag（`:8008`）；`CreateVectorWait(391 + 8·yieldable)`；post-wait dummy `SflagImmPtr(GetDummySyncFlagNumber(), "dummy sync flag")` via vtable `+1272`；conditional `CreateTracingVectorWait` gated on an `AutoProto` flag。
> - `VsyncRead` @ `0x1d524220`：`PushSyncFlagToFifo` 后接 `PopSfrFifo`（target vtable `+640`）或 `CreateScalarV2SPop`。
> - 按代常量：`GetDummySyncFlagNumber` JF=7（`0x1d4908c0`）/ PF,VF,GL=0；`SyncFlagCountModeBitOffset` 在 VF-non-lite（`0x1d49bca0`）和 GL（`0x1d4988c0`）上 = 272，其他位置为 0 — 精确立即数。
> - `mlir::llo` op classes `VSyncSetOp` / `VSyncAddOp` / `VSyncAddDoneOp` / `VSyncAddRemoteOp` / `VSyncReadOp` / `VWaitGeOp` 存在于 registered-operation tables；`num_sync_flags_encoded_per_sreg == 2` 和 `AllocationHighWaterMarkAttr::getSflagWords` @ `0x1458fec0` 确认按字粒度 packing。
>
> **[LOW]** (1) count-mode bit 272 所在的多字 SFLAG 寄存器的完整 bit-field 布局（counter 宽度、done bit、public bit）未枚举。(2) 字面按 codename 的 SFLAG 字节/字大小位于 `chip_parts.binarypb`（memfile 依赖）。(3) `HasExtraDoneBitInSyncFlags` / `SupportsYieldableOps` 的精确 `Target` vtable 索引未逐个钉住（谓词已确认按代存在；`+1272` 处 dummy-number vtable 槽位已解码）。

---

## 交叉引用

- [内存层次概览](overview.md) — SFLAG 作为 `MemorySpace::kSflag` = 6；六区域分类、render-id vs operand-tag 枚举拆分，以及此页 `4·n` 步幅馈入的 `BestFitAllocator` 矩阵
- [SMEM — 标量内存](smem-scalar-memory.md) — 相邻的标量层级；SFLAG 会与 SMEM 一起在 sequencer-local SRAM image 中被快照/传输
- [Barrier → SFLAG 编号绑定](../barrier/barrier-to-sflag-binding.md) — barrier 绑定到*哪个*编号（`[base, base+count)` 之上的保留五槽窗口）；本页是这些编号索引的基底
- [Barriers and Sync-Flags — 章节地图](../barrier/overview.md) — 驱动这些 `Vsync`/`Vwait` 原语的跨核心 barrier 子系统
- [Tree-Barrier Vsync Emitter](../barrier/tree-barrier-vsync.md) — global-barrier datapath 中 `VwaitGeSV` / `VsyncAdd` 的规范消费者
- [Remote-SFLAG Encoders](../barrier/remote-sflag-encoders.md) — `VsyncAddRemote` / `VsyncSetRemote` 如何在不同核心上编码目的端 flag；跨模式地址空间转换
- [vmem-allocator.md](vmem-allocator.md) — 共享的 `chip_parts.binarypb` 启动填充，提供数字 SFLAG 字/字节大小
- [返回索引](../index.md) — Part X — On-Chip Memory & DMA

SFLAG 同步标志内存协议 ​

摘要 ​

1. 内存模型 — 32-bit 字的扁平数组 ​

目的 ​

SflagImmPtr — 编号 → 指针 ​

硬件 count-mode bit（272） ​

3. 驱动原语 — Vsync* / Vwait* ​