SMEM 标量内存

本页所有地址都适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id md5 89edbbe81c5b328a958fe628a9f2207d）。该镜像未剥离；反混淆后的 C++ 符号名均逐字引用。其他版本会有所不同。

摘要

SMEM（scalar memory）是每个 TensorCore 的 Scalar Processing Unit (SPU) 私有的片上 SRAM tier。VMEM 是 MXU 和 VPU 读取的宽操作数暂存池，而 SMEM 是 SPU 的窄 scratch：一个扁平、按 word 粒度寻址的数组，用于支撑标量寄存器 spill，保存函数实参（parameter-pointer）表，在寄存器压力下物化循环计数器和地址计算结果，并承载 SPU 写入的每步完成描述符。它大致相当于 CPU 的 stack-plus-spill-slots 区域，但它通过专用标量 load/store opcode 显式寻址，从不通过硬件管理的 stack pointer 隐式访问。

重实现者应先抓住一个类比，然后立刻把它复杂化。SMEM 看起来像寄存器 spill 区，但它不由 HBM↔VMEM cost-balancer 管理。memory-space 分类在 MemorySpace enum 中把 SMEM 标为 kSmem = 5，概览也已经建立了所有 tier 共享的 two-stack 故事：编译期 placer（xla::jellyfish::ProgramMemoryAllocator）把 offset 冻结进 proto，每个 tier 一个 tpu::BestFitAllocator 在运行时回放它们。SMEM 与 VMEM 的决定性差异在于：MSA 从不给 SMEM 着色。FastMemorySpace() 在任何 Target 上都不会返回 kSmem（Viperfish/Ghostlite 上是 kVmem，Pufferfish 上是 kCmem，Jellyfish 上是 kHbm），所以 SMEM 是一个保留 tier，XLA 通过发出操作数 memory_space() 声明为 kSmem 的标量 load/store opcode 写入它，而不是通过 kAlternate/kDefault 的拉扯。放置由 opcode 语义驱动，而非 cost 驱动。

本页负责 SMEM 的标量模型：地址空间及其每代 sizing、byte-flat-vs-word-flat 双重性、Sld/Sst 标量 load/store opcode 家族及其寻址模式，以及把 word index 转换为 kSmem-tagged byte pointer 的 SmemWordImmPtr immediate-pointer 构造器。它不复述（不存在的）register-window 机制；该否定结论归 smem-register-window.md 所有；也不负责 SFLAG 原子协议（sflag-protocol.md）或 SPU bundle slot 编码（../isa/slot-spu-scalar.md）。

对重实现而言，契约是：

地址空间 — kSmem = 5；分配器层面 byte-flat，SPU 层面 word-flat；每个 Target 的字段布局（+0x470 size，+0x508 word size，+0x4CC log2）和 in-range 谓词。
每代 sizing — 字面 byte count 存放位置（chip_parts.binarypb）、已确认字段偏移、banking（JF 上 2 bank，PF/VF/GL 上 8 bank），以及 kSharedMemWord > kSmemWord 的更细粒度不变量。
标量 load/store 寻址 — 每代 Sld/Sst opcode 家族（ScalarLoadSmem / …Offset / …XY / ScalarStore…）、它们的三种寻址模式（absolute-imm、base-SREG+imm、base-SREG+disp-imm），以及 v7x fetch-and-add 扩展。
SmemWordImmPtr — word-index → kSmem-byte-pointer 构造器，它的 word * SmemWordSizeBytes 转换，以及它携带的 SmemWordSizeBytes == 4 断言。


MemorySpace 值	`kSmem = 5`（name table @ `0x21ce6b08`；`barna_core_smem = 9`，`sparse_core_sequencer_smem = 14`）
`SmemSizeBytes()`	`0x1d615e40` — `return (uint32)(Target + 0x470)`
`SmemWordSizeBytes()`	`0x1d617360` — `return (uint32)(Target + 0x508)`
`SmemWordSizeLog2()`	`0x1d617540` — `return (uint32)(Target + 0x4CC)`
`IsSmemByteAddressInRange(b)`	`0x1d6179a0` — `0 <= b && b < SmemSizeBytes()`
`SmemWordImmPtr(word, name)`	`0x1d516880` — word→byte ptr, `MemorySpace = kSmem(5)`
`AllocateScopedSmem(shape, name)`	`0x1d5182a0` — trampoline → `AllocateScopedMemory(…, 5u, …)`
编译期 placer	`ProgramMemoryAllocator::AllocateBytes` @ `0x1c629e40`（按 `MemorySpace` 分支）
运行时分配器	`tpu::BestFitAllocator`, `Config{base=0, end=SmemSizeBytes, align=granule=SmemWordSizeBytes}`
由 MSA 管理？	否 — `FastMemorySpace()` 在任何 Target 上都不是 `kSmem`（VF/GL 上是 `kVmem`，PF 上是 `kCmem`，JF 上是 `kHbm`）；按 opcode 语义放置
所属引擎	Scalar Processing Unit (SPU)；SREG 文件是寄存器 source/sink
置信度	CONFIRMED（字节锚定），除非某行或标注另有说明

1. 标量地址空间 {#1-the-scalar-address-space}

目的

SMEM 是 xla::jellyfish::MemorySpace::kSmem = 5。它是 overview §2 中列出的六个可寻址区域之一；MemorySpaceToString rodata 表（0x21ce6b08）中与 SMEM 相关的切片是：

text

5   smem                        ◀── SMEM scalar memory (this page)
9   barna_core_smem             ◀── BarnaCore SMEM sibling tier (PXC family only)
14  sparse_core_sequencer_smem  ◀── SC sequencer well-known-constant tier

不存在名为 SmemAllocator 的类。SMEM 使用与其他每个 tier 相同的 two-stack 架构：编译期使用 ProgramMemoryAllocator，运行时使用一个 tpu::BestFitAllocator，只通过构造它时使用的 Config 三元组区分。为什么一个 allocator 类服务所有 tier，见 overview §1 注意事项。

byte-flat / word-flat 双重性

SMEM 在两种粒度上寻址，重实现者必须把它们区分清楚：

分配器层面 byte-flat。 tpu::BestFitAllocator 发放相对于 base_offset = 0 的字节偏移（SMEM 从 sub-tile address 0 开始）。运行时 Config 是 {base_offset = 0, end = SmemSizeBytes(), alignment = SmemWordSizeBytes(), granule = SmemWordSizeBytes()}。alignment 按构造等于 granule，所以每个 SMEM 分配都是 word-aligned，BestFitAllocator ctor 的 "alignment is a power of two and alignment % granule == 0" 前置条件平凡成立。
SPU 层面 word-flat。 标量 load/store opcode 编码的是 word index，不是 byte address。转换 byte = word * SmemWordSizeBytes 在 LLO lowering 时应用（见 §4）。在 LLO IR 内部，值携带 StrongInt<SmemOffset>，也就是word index，所以分配器在 IR 中从不观察原始 byte address；它看到的是通过 SmemWordSizeBytes() 转换后的 word offset。strong-int arc-state 类型逐字出现在 LloCodeStepper::ArcState<SmemOffset, SregValue>::SetValue（0x10bc7ac0）中。

GOTCHA — 天真的重实现如果把 opcode immediate 当作 byte address，会在 word_index 字节处索引 SMEM，而不是 word_index * SmemWordSizeBytes 字节处，即相差一个 word size 因子。opcode immediate 永远是 word index；byte 转换是 allocator/lowering 的工作，并且只在 SmemWordImmPtr 中应用一次。

`Target` 字段布局

SMEM geometry 位于每个 codename 的 Target 对象中，由启动时的 chip_parts.binarypb 填充。直接 accessor 及其反编译 body：

// 0x1d615e40  Target::SmemSizeBytes() const
//   return *(uint32_t*)(this + 0x470);          // total SMEM bytes, signed-compared
// 0x1d617360  Target::SmemWordSizeBytes() const
//   return *(uint32_t*)(this + 0x508);          // per-bank word width, bytes (= granule)
// 0x1d617540  Target::SmemWordSizeLog2() const
//   return *(uint32_t*)(this + 0x4CC);          // == log2(SmemWordSizeBytes)
// 0x1d6179a0  Target::IsSmemByteAddressInRange(int b) const
//   return b >= 0 && b < *(int*)(this + 0x470); // 0 <= b < SmemSizeBytes()
// 0x1d618140  Target::SmemUserSpaceWordOffset() const
//   return *(uint64_t*)(this + 0x7F8);          // first user-allocatable word index

字段	偏移	类型	含义
`SmemSizeBytes`	`+0x470`	int32	SMEM 总字节数（按 signed 比较；in-range check 使用 `>= 0`）
`BarnaCoreSmemSizeBytes`	`+0x47C`	int32	BarnaCore SMEM sibling-tier size
`BarnaCoreSmemBaseBytes`	`+0x480`	int32	BarnaCore window 在 SMEM 内部开始处的 byte offset
`SmemWordSizeLog2`	`+0x4CC`	uint32	`log2(SmemWordSizeBytes)`，opcode encoder 使用的 shift
`SmemWordSizeBytes`	`+0x508`	uint32	每 bank 的 word 宽度；alignment 和 granule 都等于它
`BarnaCoreSmemWordSizeBytes`	`+0x51C`	uint32	BarnaCore SMEM word width
`SmemUserSpaceWordOffset`	`+0x7F8`	uint64	用户分配可使用的第一个 word index（高于底部 reservation）
`BarnaCoreFregSmemWordOffset`	`+0x818`	uint64	BarnaCore freg SMEM word offset

NOTE — SmemSizeBytes 按 int32 读取，range check 是 b >= 0 && b < SmemSizeBytes()（IsSmemByteAddressInRange, 0x1d6179a0）。重实现必须保留下界 >= 0 guard：该字段是 signed，负 byte address 与"太大"是不同的拒绝情形。匹配的编译期诊断是 "byte_address < target().SmemSizeBytes()"。

派生偏移

两个 accessor 从 param-pointer table 而不是存储字段计算保留区域边界：

// 0x1d618180  Target::ReservedSmemInBytes(int i) const
//   return SmemSizeBytes()
//        - (ParamPtrLocationWordOffset(i) - 1) * SmemWordSizeBytes();
// 0x1d618040  Target::StartReservedSmemWordOffset(int i) const
//   tail-call → ParamPtrLocationWordOffset(i)   (0x1d617fa0)

StartReservedSmemWordOffset(i) 是对 ParamPtrLocationWordOffset(i) 的字面 tail call，因此 parameter-pointer table 和 SMEM 顶部 reserved-block table 共享相同的 offset arithmetic：reserved block 位于 SMEM 顶部，紧贴 parameter pointer 放置位置之下。ReservedSmemInBytes(i) 随后把该 word 边界转换为 image 顶部 reserved span 的 byte count。

NOTE — ParamPtrLocationWordOffset(int)（0x1d617fa0）的字面内容，也就是驱动两个派生 accessor 的每个 i 常量表，尚未逐项追踪。上面两个 accessor 已确认；它们索引的表在反编译前为 LOW confidence。

2. 每代 Sizing {#2-per-generation-sizing}

目的

§1 中的 accessor 读取四个 Target 字段；字面值由每个 codename 的 TpuChipParts::FromProto / TpuMemoryParts::FromProto 从嵌入的 chip_parts.binarypb 放入。C++ accessor、banking 和 capability flag 已完全解码；数值 byte size 仍位于尚未解码的 protobuf 中。

字节锚定的 sizing 事实

Generation	Target class	SMEM banks	`ScalarLoadLatency` (cy)	4-byte SMEM-write DMA opcode	SCS SMEM (bytes)
v2 / Jellyfish (JF)	`JellyfishTarget`	2	2	`false` (`0x1d48fee0`)	n/a (no SCS)
v4 / Pufferfish (PF)	`PufferfishTarget`	8	4	`false` (`0x1d4946a0`)	0 (no SCS in SMEM; `0x1fbac3e0`)
v5 / Viperfish (VF)	`ViperfishTarget`	8	6	`true` (`0x1d49a8e0`)	64 KiB (vfc) / 0 (vlc)
v6e / Ghostlite (GL)	`GhostliteTarget`	8	6	`true` (`0x1d4976c0`)	64 KiB (`0x1fe6dd60`)
v7x / `6acc60406` (gfc)	`GhostliteTarget`	8	6	`true`	64 KiB (`0x1fda8aa0`)

Banks 来自每个 Target 的 MemBanks(MemorySpace) override 在 kSmem(5) 分支上的返回值：JellyfishTarget::MemBanks（0x1d48fc80）对 kSmem 返回 2；PF/VF/GL 返回 8。byte offset B 的 bank index 是 (B / SmemWordSizeBytes) mod MemBanks(kSmem)；bank 内 row 是 (B / SmemWordSizeBytes) / MemBanks(kSmem)（按与 VMEM 的对称性假定，精确公式为 MEDIUM confidence；bank count 为 CERTAIN）。
ScalarLoadLatency（load-to-use cycles，SREG ← SMEM）是 LLO bundle packer 遵守的硬延迟：JF=2，PF=4，VF=6，GL=6。
Supports4BSmemWriteDmaDestinationOpcode() 是硬 capability flag。VF/GL 暴露 coalesced 4-byte SMEM-write DMA termination opcode（启用 "AllToAllDynamic with smem4b mode enabled." codepath）；JF/PF 必须使用更宽的 32-byte DMA wires（"Pipelined AllToAllDynamic is supported for only non-smem4b mode"）。Target:: base 是 pure-virtual LogMessageFatal（0x1d61d500）。
SCS SMEM 是 SparseCore Scalar-Smem 子集，由 asic_sw::deepsea::<family>::HardwareAttributes::GetSparseCoreScsSmemSizeBytes 直接返回编译期常量；gfc/glc 的反编译 body 字面上是 return 0x10000;（64 KiB）。viperfish-lite shim（vxc::vlc, 0x1ff4b7a0）和 Pufferfish（0x1fbac3e0）返回 0。

QUIRK — SCS SMEM 大小是二进制中的硬编码 immediate（mov $0x10000, %eax），不是 chip_parts.binarypb 字段；这与通用 SmemSizeBytes 不同。因此 SparseCore scalar-memory 预算在 Viperfish/Ghostlite/Ghostfish 之间相同（64 KiB），不能通过替换 protobuf 重新定目标；它按每个 HardwareAttributes subclass 编译进来。

两个更细粒度不变量

SMEM 是两个片上标量 tier 中粒度更细的那个。两个 assertion 强制这一点；每个违反时都是硬 LogMessageFatal：

text

"kSharedMemWordSizeBytes > kSmemWordSizeBytes"
"kSharedMemWordSizeBytes % kSmemWordSizeBytes == 0"
"hbm_word_size_bytes > smem_word_size_bytes"
"hbm_word_size_bytes % smem_word_size_bytes == 0"

HBM 和 shared-memory word 是 SMEM word 的严格倍数，并且严格更大。选择 SMEM word ≥ HBM word 的重实现会在 lowering 时触发 LogMessageFatal。

尚未解码的内容

NOTE — 每个 codename 的字面 SmemSizeBytes 和 SmemWordSizeBytes 数值位于 chip_parts.binarypb 中，尚未提取（数值为 LOW confidence）。上面的字段偏移和 accessor body 为 CERTAIN。有一个具体数字已钉住：SmemWordImmPtr（§4）断言对其编译目标 codename，SmemWordSizeBytes() == sizeof(uint32_t)，所以该路径上的 word 是 4 bytes。公开 TPU SPU-scratch 预算暗示总 SMEM 在 {16, 64, 128, 256} KiB 中，但字面数值尚未确认。

BarnaCore SMEM 兄弟 tier

BarnaCoreSmem（barna_core_smem = 9）是第二个 SMEM tier，有自己的 size/base/word-size 字段（+0x47C、+0x480、+0x51C）以及由 SupportsBarnaCore()（vtable[+0x258]）gate 的 accessor；否则 BarnaCoreSmemSizeBytes 会 LogMessageFatal "BarnaCore is not supported by this target"。该二进制中只有 Pufferfish family 带有非空 BarnaCore window；其标量 load/store opcode 使用 BarnaCoreSequencerScalar1_ 前缀，而不是 TensorCoreScalar1_。匹配的 scoped-frame trampoline 是 AllocateScopedBarnaCoreSmem（0x1d518500，MemorySpace = kBarnaCoreSmem）。

3. 标量 Load / Store 寻址（`Sld` / `Sst`）

目的

SPU 通过专用标量 load/store opcode 访问 SMEM，即 Sld/Sst 家族。这些不是 DMA，cost model 也不会把它们建模为 DMA；它们经由专用 SREG-read port 退出，整个 round trip 由 ScalarLoadLatency cycles 覆盖（§2）。这些 opcode 的 bundle-slot 编码归 ../isa/slot-spu-scalar.md 所有；本节记录寻址模型，即每个 opcode 计算什么 SMEM 地址。

三种寻址模式

不同 generation 的 opcode 名称不同，但寻址模式可归并为三种。opcode-name decode 已根据 *_functions.json 符号表确认（ScalarLoadSmem、ScalarLoadSmemOffset、ScalarLoadSmemXY 家族，每个都有 …AddressField / …OffsetField / …DestField encoder accessor）。

模式	计算出的地址	PXC opcode (TPU v4)	VXC/GXC opcode (v5/v5+)
Absolute immediate	`SMEM[imm_word]`	`ScalarLoadSmem`	`ScalarLoadSmemY`
Base-SREG + imm	`SMEM[SREG_base + imm]`	`ScalarLoadSmemOffset`	—
Base-SREG + disp-imm	`SMEM[SREG_x + imm_disp]`	—	`ScalarLoadSmemXY`
Store, absolute	`SMEM[imm_word] := SREG`	`ScalarStoreSmemAbsolute`	`ScalarStoreXToSmemY`

PXC / Scalar1 slot（Pufferfish，也包括 Jellyfish 形态）：
- TensorCoreScalar1_ScalarLoadSmem — 在 immediate-encoded absolute word address 处 load 一个 SMEM word 到 SREG。immediate 通过 Place16BitImmediate / Place32BitScalarImmediate 放置；encoder 转换 word_offset = byte_offset / SmemWordSizeBytes。
- TensorCoreScalar1_ScalarLoadSmemOffset — 同样的 load，地址 = base_SREG + immediate。这是 parameter-table / stack-frame 访问模式，其中 base 位于 SREG 中，displacement 是常量。
- TensorCoreScalar1_ScalarStoreSmemAbsolute — 在 immediate absolute address 处执行 SREG → SMEM。用于 sync-flag write、completion-descriptor write、return-value write。
VXC/GXC / ScalarAlu slot（Viperfish，Ghostlite）：
- TensorCoreScalarAlu_ScalarLoadSmemY — SREG_dest ← SMEM[Y]，Y 是 16-bit immediate word index。
- TensorCoreScalarAlu_ScalarLoadSmemXY — SREG_dest ← SMEM[X + Y]，X 是 SREG base，Y 是 16-bit immediate displacement。这是 parameter-table read 使用的 canonical "load with displacement"。
- TensorCoreScalarAlu_ScalarStoreXToSmemY — SMEM[Y] ← SREG_X。
- SparseCore SPU 变体存在，前缀为 SparseCoreScalarAlu_，目标是 SCS SMEM 子空间。

QUIRK — PXC family 和 VXC/GXC family 不是彼此的重命名。PXC 有单个 …Offset 形式（base + imm）；VXC/GXC 拆成 …Y（absolute imm）和 …XY（base + disp）。把 ScalarLoadSmemOffset 直接映射到 ScalarLoadSmemXY 的重实现会在其中一代上弄错操作数角色：PXC 上 SREG 是 base；VXC 上 SREG（X）是 base，immediate（Y）是 displacement，但 VXC 上 absolute form 是单独的 …Y opcode，完全没有 SREG 操作数。

v7x fetch-and-add 扩展

Ghostfish（GXC/GFC）增加了 JXC/PXC 中不存在的两个 atomic-reduce-into-SMEM opcode：

SparseCoreScalarAlu_ScalarStoreXToSmemSumDestAndY — atomic SMEM[Y] := SMEM[Y] + SREG_X。这是 SMEM 层面的单指令 fetch-add，用在多个 SPU 向共享 SMEM word 归约的位置。
SparseCoreScalarMisc_SmemFetchAndAdd — 正交 ScalarMisc pipe 中的同一操作（因此可与普通 scalar load 并发 issue）；把 pre-add SMEM 值返回到 destination SREG。由 isa_emitter::EmitFetchAndAddOp<glc::SparseCoreTecBundle, SmemFetchAndAdd>（0x13a3a300）发出。存在于 VFC 和 GLC/GFC；不存在于 JXC/PXC。

落入 SMEM 的内容

SMEM 保存 SPU 的标量工作集：

标量寄存器 spill — 当 LSRA-v2 无法让某个 SREG 在一个区域内保持 live 时，它会发出 ScalarStore…ToSmem… 到保留 spill 区域，然后在下一次 use 处发出 ScalarLoadSmem…。SMEM 是 SREG 文件的 spill backing store，不是它之上的窗口（见 §5）。
parameter-pointer table — ParamPtrLocationWordOffset(int)（0x1d617fa0）给出 parameter-pointer table 所在的 SMEM word。LLO ScalarLoadSmem lowering 会读取它，在每个 kernel 开头把函数实参物化到 SREG 中。
SC-sequencer well-known constants — chip_id、replica_id、partition_id、subslice_origin、hbm_offset。LloAddress::MakeSparseCoreSequencerSmemConstant(long)（0x1d60bc60）在每个 codename 的硬编码 SMEM offset 处构造一个 LloAddress；GetIntegerFromSmemOpLowering MLIR pattern 会把每一个 lowering 为该固定地址的 ScalarLoadSmem，完全绕过 allocator。
保留的顶部 word — last-set P/T-state（"Reserved Smem offset for storing last set P/T-state."）、trace context（TpuChipProfilerVxcImpl::ReadTraceContextFromSmem, 0x1d1a22a0）、completion descriptor。
保留的底部 word — SCS overlay trampoline、ProgramContinuator stack frame（GetTensorCoreStackSizeInWordForSparseCoreSmem 0x1d17cb60，GetSparseCoreStackSizeInWordForSparseCoreSmem 0x1d17cc80）。

GOTCHA — SMEM 不会隐式 zero-initialized。buffer-assignment pre-pass 明确拒绝 zero 一个 SMEM allocation："Cannot zero out AllocateBuffer output memory space is smem"。分配 SMEM 区域的一方必须在读取前写入它依赖的每个 word；假设 scratch 零填充的重实现会读到陈旧数据。

MLIR / LLVM dialect ops

Sld/Sst opcode 的编译器侧等价物：

text

llo.alloca_smem        (mlir::llo::AllocaSmemOp)        — SMEM stack-frame allocation; getNumWords()
llo.saddr.smem         (mlir::llo::ScalarAddressSmemOp) — materialize an SMEM address into an SREG
                          (read effect on mlir::sparse_core::resource_effects::Smem)
llvm.tpu.alloca.smem / llvm.tpu.allocate.smem / .allocate.smem.any   — LLVM intrinsic equivalents
llvm.tpu.addrspacecast.smem                              — bridge LLO SMEM addr-space ↔ generic LLVM ptr
llvm.tpu.dma.hbm.to.smem.sc.{simple,general,single.strided}  — SparseCore HBM→SMEM DMA variants

NOTE — 数值 LlvmTpuDialect::SmemAddressSpace() 由 sentinel "address_space == LlvmTpuDialect::SmemAddressSpace()" 断言。通用 SMEM LLVM address space 是 0；MemorySpace 1（smem）在 MemorySpaceToAddressSpace 反向表中映射到 LLVM addrspace 0（dword_AF36CE8[0] == 0），已在 Address-Space ID Table 上确认。LLO kSmem MemorySpace enum 值（5）与这个 LLVM address space 属于不同编号空间，不能混淆。

4. `SmemWordImmPtr` — Immediate-Pointer 构造器 {#4-smemwordimmptr--the-immediate-pointer-constructor}

目的

LloRegionBuilder::SmemWordImmPtr(long word_index, std::string_view name)（0x1d516880）是把 word index 转换为 LLO IR 中 kSmem-tagged byte pointer 的唯一 chokepoint。每个通过常量 word offset 寻址 SMEM 的标量 load/store 都经过它。这也是 §1 中 byte-flat/word-flat 双重性被解析的位置：恰好一次，就在这里。

算法

反编译 body（清理掉 absl::StatusOr 机械结构）如下：

// 0x1d516880  LloRegionBuilder::SmemWordImmPtr(long word_index, string_view name)
LloValue* SmemWordImmPtr(LloRegionBuilder* rb, long word_index, string_view name):
    Target* tgt = rb->module->target;                 // *(*(rb)+56)+16

    // Invariant: the SMEM word must be exactly a uint32 for this path.
    word_bytes = tgt->SmemWordSizeBytes();             // 0x1d617360
    CHECK_EQ(word_bytes, sizeof(uint32_t)):            // "target().SmemWordSizeBytes() == sizeof(uint32_t)"
        // on failure → LloModule::UpdateStatus(... CheckFailer ...)
        //   site: platforms/xla/service/jellyfish/llo_region_builder.cc

    word_bytes = tgt->SmemWordSizeBytes();             // re-read after the check
    Shape shape = ShapeUtil::MakeValidatedShape(U32 /*=8*/, /*rank*/0);  // a 4-byte scalar

    // Convert word → byte and build the immediate pointer, tagged kSmem.
    return rb->ImmPtr(/*byte_offset=*/ word_bytes * word_index,   // ◀── the word→byte conversion
                      shape,
                      /*MemorySpace=*/ 5 /*kSmem*/,               // ◀── tier tag
                      name, ...);

重实现者必须精确复现三件事：

word→byte 转换是 SmemWordSizeBytes() * word_index，即乘法，在 SmemWordImmPtr 内应用一次。调用者传入 word index；IR pointer 携带 byte offset。
元素类型是 4-byte unsigned scalar。 MakeValidatedShape(8, …) 构造 U32 rank-0 shape；immediate pointer 指向单个 32-bit SMEM word。
传给 ImmPtr 的 MemorySpace 参数是字面 5（kSmem）。 这声明结果 pointer 属于 SMEM tier，进而使 ProgramMemoryAllocator::AllocateBytes 把值提交到 SMEM image，而不是 VMEM/HBM。放置由 opcode/pointer 语义驱动，正如 § 摘要和概览所述。

QUIRK — SmemWordImmPtr 硬断言 SmemWordSizeBytes() == sizeof(uint32_t)。Target 字段是可容纳任意 2 的幂的 uint32，但这个 immediate-pointer 路径只支持 4-byte SMEM word。如果某个 codename 的 chip_parts.binarypb 设置了非 4-byte SMEM word，那么一旦构造任何 constant-offset SMEM pointer 就会 LogMessageFatal。因此对 0.0.40 的生产 codename 而言，该路径上的 word 实际上被固定为 4 bytes。

Scoped-frame 分配

对每区域 SMEM scratch（不是固定 word 常量），入口是 trampoline AllocateScopedSmem：

// 0x1d5182a0  LloRegionBuilder::AllocateScopedSmem(Shape const&, string_view name)
LloValue* AllocateScopedSmem(rb, shape, name):
    return rb->AllocateScopedMemory(shape, /*MemorySpace=*/ 5u /*kSmem*/, name);  // 0x1d517c20

反编译显示它是单个 tail call，传递字面 5u，即 scoped-VMEM 的 SMEM 类比，只在 MemorySpace 参数上不同（5 而不是 3）。AllocateScopedMemory 委托给 LloRegion::AllocateScopedFixedMemory（0x1d5137c0）。BarnaCore 兄弟 AllocateScopedBarnaCoreSmem（0x1d518500）传递 BarnaCore MemorySpace。

LSRA-v2 spill-window 算术

LSRA-v2 register allocator 使用同一组 accessor 计算可用于 spilling 的 SMEM 字节数。lsrav2::SmemBytesAvailable(LloRegion*)（0x12786120）反编译为：

// 0x12786120  lsrav2::SmemBytesAvailable(const LloRegion*)
long SmemBytesAvailable(self, region):
    Target* tgt = ...;
    int reservation_id = ...;                          // region field +52
    word_bytes  = tgt->SmemWordSizeBytes();
    ceiling     = tgt->ParamPtrLocationWordOffset(0);  // top user word (param-ptr table)
    floor       = tgt->SmemUserSpaceWordOffset();       // first user word (above bottom blocks)
    span_bytes  = word_bytes * (ceiling - floor);       // raw user window in bytes
    hbm_reserve = round_up(tgt->HbmWordSizeBytes(), 1024);     // HBM-aligned head reserve
    return span_bytes - (hbm_reserve + tgt->ReservedSmemInBytes(reservation_id));

spill window 是用户 span（从 SmemUserSpaceWordOffset floor 到 ParamPtrLocationWordOffset(0) ceiling）减去一个 HBM-word-aligned reserve 和 top-of-SMEM reserved span。这确认 SmemUserSpaceWordOffset（+0x7F8）的角色是用户可分配区域的底部，ParamPtrLocationWordOffset(0) 是其顶部。

NOTE — spill-region cap 还可由 FLAGS_xla_jf_lsra_v2_reserved_smem（0x223afaa8）调节；该 flag 保留固定 top-N words 作为 scratch，并把 cap 以上的 load 标记为 rematerializable（"Considers all smem loads above the spill limit to be const and read-only and really trivially rematerializable."）。该 flag 到 assertion "current_local_sync_flag_ <= first_smem_scratch_word_" 所引用的 first_smem_scratch_word_ 字段的精确翻译尚未追踪（该特定 wiring 为 LOW confidence）。

5. 没有寄存器窗口 {#5-no-register-window}

SMEM 没有 register-window 机制。对 SmemRegisterWindow / SmemRegisterFile / SregWindow / SmemSpillRegister 的二进制搜索返回 zero hits。SMEM 是扁平 byte/word array，不是寄存器文件上的窗口；它没有 SPARC 风格的 register-window overflow 故事。标量 register windowing 完全位于 SREG file（由 LSRA-v2 驱动的 xla::jellyfish::SregNumber typed pool），SMEM 只是该文件的 spill backing store。完整否定结论和 SREG-file 细节归 smem-register-window.md 所有；本页只记录 SMEM 标量模型按设计无窗口。

6. Exhaustion 和不变量处理

SMEM overflow 几乎完全是编译期问题，因为 MSA 不会 rebalance SMEM，image 在执行前已经完全 layout（运行时 allocator 只是回放）。编译期路径：

模式	触发条件	诊断
Out-of-range byte offset	`IsSmemByteAddressInRange(b)` false (`0x1d6179a0`)	`"byte_address < target().SmemSizeBytes()"`
Allocator cannot place	shared `BestFitAllocator::Allocate` OOM	`absl::ResourceExhaustedError(...)`（见 hbm-allocator.md）
Bad `smem_end`/`smem_start` clamp	control-plane config reject	`"Can't set smem_end using too large of a value."`
Geometry invariant	word-size / granule violations	`"kSharedMemWordSizeBytes > kSmemWordSizeBytes"`, `"hbm_word_size_bytes % smem_word_size_bytes == 0"`, `"available_smem_size >= granule_size"`
SC table overflow	ragged-pointer table exceeds SCS budget	`"Row pointers would exceed available SCS Smem ("`
Trampoline overflow	SCS overlay trampoline reservation too large	`"Reserve extra smem spill area for SCS overlays trampoline."` (`FLAGS_xla_sc_reserve_scs_trampoline_smem`, `0x22335e88`)
Debug poison	use-after-free detector	`"Poisoned Smem value use detected"` (`xla::jellyfish::llo_analysis::RaceAnalyzerStepper::PoisonSmemBuffer`, `0x10bc15c0`)

每个 geometry-invariant message 都是硬 LogMessageFatal。SCS-budget check（GetUserAllocatableScsSmemSize, 0x13db6d80）报告编译期错误，而不是静默 spilling，因此溢出 SCS SMEM 的 SparseCore lowering 会响亮失败。

7. 编译期 → 运行时交接

SMEM 遵循共享 hand-off pipeline（overview §3，hbm-allocator.md），只在 MSA 阶段分歧：

text

1. HeapSimulator::Run(GlobalDecreasingSizeBestFitHeap, budget = SmemSizeBytes())
2. MSA pass — DOES NOT relocate SMEM (SMEM is not kAlternate; FastMemorySpace() is never kSmem — kVmem on VF/GL, kCmem on PF, kHbm on JF)
3. ProgramMemoryAllocator::AllocateHloBuffer (0x1c62a5a0)
     emits ProgramMemoryMetadata_Allocation{ memory_space = kSmem, offset, size, block_type, name }
4. proto travels inside the compiled XDB / LLO program
5. ProgramMemoryAllocator::CreateFromProto (0x1c631f20) rehydrates runtime state
6. TpuHal binds one tpu::BestFitAllocator for the SMEM tier:
     Config{ base_offset = 0,
             end        = Target::SmemSizeBytes(),
             alignment  = Target::SmemWordSizeBytes(),
             granule    = Target::SmemWordSizeBytes() }

决定性分歧是阶段 2：SMEM 从不由 MSA 着色。某个值落在 SMEM 中，是因为它的 lowering 发出了声明 memory_space() == kSmem 的标量 load/store opcode（在每个 emission site 由 "address->memory_space() == MemorySpace::kSmem" 和 "dest_address->memory_space() == MemorySpace::kSmem" 断言），然后 AllocateBytes(MemorySpace = kSmem, …) 把它提交进 SMEM image。SMEM 没有 cost-balancing tug-of-war。

名称	关系
内存层级概览	负责 six-region taxonomy 和 `MemorySpace` enum；SMEM 是 `kSmem = 5`
SMEM 寄存器窗口	否定结论：SMEM 没有 register window；SREG-file windowing 细节
SFLAG 协议	兄弟 `kSflag = 6` tier；其 size/word 字段紧邻 SMEM 字段（`+0x468`, `+0x504`）
CMEM 池	兄弟片上 operand pool（`kCmem`，仅 Pufferfish）；与 SMEM 不同，受 MSA 管理
HBM 分配器	SMEM 在运行时回放的 `BestFitAllocator` 算法；共享 OOM 路径
VMEM 分配器	与 SMEM 对比的 `kAlternate` MSA-managed tier

交叉引用

overview.md — memory-space taxonomy、enum、two-stack allocator story；SMEM tier 行
smem-register-window.md — 为什么 "register window" 不适用于 SMEM；SREG-file spill 模型
sflag-protocol.md — SPU 也会访问的独立 kSflag atomic tier
cmem-pool.md — Pufferfish CMEM operand pool，MSA-managed sibling
hbm-allocator.md — tpu::BestFitAllocator 算法、Config 三元组、ResourceExhaustedError
../isa/slot-spu-scalar.md — Sld/Sst opcode 的 SPU bundle-slot 编码
../isa/memory-space-enum.md — 用作 SMEM tier tag 的 MemorySpace enum 值

SMEM 标量内存 ​

摘要 ​

1. 标量地址空间 {#1-the-scalar-address-space} ​

目的 ​

byte-flat / word-flat 双重性 ​

Target 字段布局 ​

派生偏移 ​

2. 每代 Sizing {#2-per-generation-sizing} ​

目的 ​

字节锚定的 sizing 事实 ​

两个更细粒度不变量 ​

尚未解码的内容 ​

BarnaCore SMEM 兄弟 tier ​

3. 标量 Load / Store 寻址（Sld / Sst） ​

目的 ​

三种寻址模式 ​

v7x fetch-and-add 扩展 ​

落入 SMEM 的内容 ​

MLIR / LLVM dialect ops ​

4. SmemWordImmPtr — Immediate-Pointer 构造器 {#4-smemwordimmptr--the-immediate-pointer-constructor} ​

目的 ​

算法 ​

Scoped-frame 分配 ​

LSRA-v2 spill-window 算术 ​

5. 没有寄存器窗口 {#5-no-register-window} ​

6. Exhaustion 和不变量处理 ​

7. 编译期 → 运行时交接 ​

相关组件 ​

交叉引用 ​

SMEM 标量内存

摘要

1. 标量地址空间 {#1-the-scalar-address-space}

目的

byte-flat / word-flat 双重性

`Target` 字段布局

派生偏移

2. 每代 Sizing {#2-per-generation-sizing}

目的

字节锚定的 sizing 事实

两个更细粒度不变量

尚未解码的内容

BarnaCore SMEM 兄弟 tier

3. 标量 Load / Store 寻址（`Sld` / `Sst`）

目的

三种寻址模式

v7x fetch-and-add 扩展

落入 SMEM 的内容

MLIR / LLVM dialect ops

4. `SmemWordImmPtr` — Immediate-Pointer 构造器 {#4-smemwordimmptr--the-immediate-pointer-constructor}

目的

算法

Scoped-frame 分配

LSRA-v2 spill-window 算术

5. 没有寄存器窗口 {#5-no-register-window}

6. Exhaustion 和不变量处理

7. 编译期 → 运行时交接

相关组件

交叉引用