DeepInfer 源码拆解（十八）：Engram 条件记忆的推理路径——词表压缩、N-gram 哈希与跨 rank 查表

一个容易被忽略的约束：只支持 PP 放置

v4_modules/engram_parameters.py 里有一行很短的断言：

1
assert engram_cfg.placement == "pp"
2
assert engram_cfg.fp8_proj or not engram_cfg.fp8_embed

placement = "pp" 意味着 Engram 层固定挂在主干流水的位置上，不能像某些并行方案那样把整层挪到单独的 rank 上做「记忆服务器」。原因不复杂：Engram 的输出要写回主干的 hidden states，而 V4.1-Flash 用的是 mHC 多路残差主干，写入必须发生在持有该层主干的 rank 上。把表切到更多 rank 是 ENP 的事，把计算挪走则是另一回事。

第二条断言约束了量化组合：如果嵌入表用 FP8 存储，投影矩阵也必须用 FP8 GEMM。否则取回的 FP8 行要先反量化成 BF16 再做 BF16 GEMM，白白多一次转换；这条规则把「省显存」和「省带宽」绑在一起。

Engram 层不能落在 YOCO 的 consumer 半区

models/deepseek.py 的 _init_yoco 里还有一条断言：

1
assert all(layer_id < cfg.n_prefill_layer for layer_id in engram_layer_ids), (
2
    "yoco consumer compaction does not support Engram in consumer layers yet."
3
)

YOCO 在 prefill 阶段会把序列在某个边界层截断，consumer 半区的行布局是 compact 的，而 ngram_token_ids、images_seq_mask 都是按 full 布局准备的。如果 Engram 落在 consumer 层，查表用的 token 行号与 hidden states 行号就对不上。当前配置把两个 Engram 层都放在边界之前，这条断言只是把「暂不支持」写成了运行时检查。读到这里可以记住一个判断方法：凡是按行索引 hidden states 的模块，都要检查它是否兼容 YOCO 的 compact 视图。

为什么需要一个「压缩词表」

Engram 的键是 token 序列本身。如果直接用原始 token id 做哈希，会碰到两个问题。

第一，同一个词在原始词表中可能有多个 id。Tokenizer 的训练会把高频词切成整词、把低频词切成分段，同一个字符串在分词后是哪个 id 取决于上下文和空白前缀。例如「 hello」与「hello」在很多 tokenizer 里是两个不同的 id；大小写变体、全角与半角、带重音字母的不同写法，也可能各自占一行。如果直接拿 id 哈希，语义相同的 N-gram 会落到完全不同的表项上，表的容量被这些变体摊薄，模型要靠训练去学「这两个 id 其实是同一个词」，效率很低。

第二，4 亿行的表里，如果大量行只对应极少出现的形式，等价于浪费。论文说 Engram 的检索是确定性的，那就应该在检索前把显而易见等价的 token 合并掉，让有限的行数花在真正不同的语义单元上。

于是 DeepInfer 在推理时对每个 token 做了一次「原始 id → 压缩 id」的映射。这个映射由 CompressedVocabulary 在初始化时构建，运行时的查询由 CompressedTokenizer 完成。

压缩映射怎么构造

engram_metadata.py 中的 CompressedVocabulary.__init__ 很短，但每一步都有目的：

1
normalizer = normalizers.Sequence([
2
    normalizers.NFKC(),
3
    normalizers.NFD(),
4
    normalizers.StripAccents(),
5
    normalizers.Lowercase(),
6
    normalizers.Replace(Regex(r"[ \t\r\n]+"), " "),
7
    normalizers.Replace(Regex(r"^ $"), sentinel),
8
    normalizers.Strip(),
9
    normalizers.Replace(sentinel, " "),
10
])
11
key_to_id: dict[str, int] = {}
12
self.old2new: list[int] = []
13
for token_id in range(self.tokenizer.get_vocab_size()):
14
    text = self.tokenizer.decode([token_id], skip_special_tokens=False)
15
    if "�" in text:
16
        key = self.tokenizer.id_to_token(token_id)
17
        assert key is not None
18
    else:
19
        key = normalizer.normalize_str(text) or text
20
    self.old2new.append(key_to_id.setdefault(key, len(key_to_id)))

逐条解释：

• NFKC 把兼容字符折叠成标准形式，全角字母、全角数字、上标等会与半角形式合并。
• NFD + StripAccents 把带重音的字母拆成「基字母 + 组合符」再去掉组合符，é 与 e 落到同一个键上。这对多语言模型尤其重要：同一个词在不同语言的文本里有不同拼法，底层语义单元是字母本身。
• Lowercase 合并大小写。注意这一步对中文没有影响，但对英文与部分拉丁语系词汇的收益很大。
• 空白折叠：所有连续空白（空格、制表、换行）折叠成一个空格。这是为了处理 tokenizer 中「前导空格」这类形式：▁hello、 hello、\nhello 归一化后如果都变成同一个键，就合并进同一行。
• 前后 Strip + 哨兵替换：先把一个纯空格的键换成  哨兵，执行 Strip，再把哨兵换回空格。这样做的原因是 Strip 会把纯空白键变成空串，而空串和「没有内容」无法区分；用哨兵占位可以保证「空白 token」单独拥有一个键，不与其它 token 共享。
• 非法 UTF-8 的处理：tokenizer.decode 出来的文本里如果含替换字符 �，说明这个 token 是无法单独解码的字节片段（byte-level BPE 的中间片段）。这类 token 不能走规范化，否则不同字节片段会被折叠成同一个 � 串而错误合并；此时直接用词表里的原始字符串 id_to_token 作为键。
• setdefault 保证确定性编号：按 token id 升序遍历，第一次遇到某个键时分配新压缩 id，之后遇到的相同键复用同一 id。由于遍历顺序固定、字典插入顺序固定，不同进程、不同机器上构建出的映射完全一致。这一点在 EP 场景下是必须的：每个 rank 都要对同一批 token 算出同样的哈希，才能定位到同一行。

self.lookup_table = np.asarray(self.old2new, dtype=np.int32) 是最终的映射表，长度等于原始词表大小，下标是原始 id，值是压缩 id。

运行时：一次整批 token 的映射

压缩映射在推理时不是逐 token 逐个查字典，而是一次性用一个 C++ 单遍函数完成：

1
@classmethod
2
def encode(cls, token_ids: np.ndarray[np.int32]) -> np.ndarray[np.int32]:
3
    assert cls.lookup_table_np is not None, "CompressedTokenizer is not initialized"
4
    return encode_compressed(token_ids, cls.lookup_table_np, cls.IMAGE_SENTINEL)

encode_compressed 来自 deepinfer._C.engine，做的是「对数组每个元素做 lookup_table[id]，越界则填哨兵」。为什么放在 C++ 里？因为输入序列可能是几万到几十万 token（prefill），Python 循环或 numpy fancy indexing 都会带来一次额外的中间数组分配；单遍映射可以原地读、原地写，把这一步压到可以忽略。CompressedTokenizer 还同时保存了 lookup_table_gpu，供 GPU 侧 kernel 做同样的映射（后文讲投机解码时会用到）。

IMAGE_SENTINEL = -1 是图像 token 在压缩空间的表示。图像 token 不对应任何文本词，也就没有压缩 id；把它映射成 -1 而不是某个合法 id，是为了让下游能够区分「这是图像」与「这是某个文本 token」。

压缩之后，词表还剩多大

CompressedVocabulary.num_tokens 记录了压缩后的词表大小，NgramHashMapping 用它作为哈希取模的上界。这个大小取决于 tokenizer，无法给出固定值，但可以确定的是：压缩后的词表是原始词表的一个「语义化」子集，NFKC、去重音、小写与空白折叠共同去掉了大量近似重复项。工程上要关心的不是压缩比，而是压缩映射必须在所有 rank 上一致——否则同一句话在不同 rank 上算出不同的压缩 id，落在不同的表行，取回的 embedding 就会错位。

一个隐藏的初始化时机

CompressedTokenizer.initialize(tokenizer_path) 是一个类级单例初始化，构建 CompressedVocabulary、保存 CPU 侧 numpy 映射表、把映射表上传到 GPU，并把 pad token 替换为压缩后的 id。release() 会把这些清空。源码注释写着「tokenizer 压缩关闭时 initialize 不调用」，说明这条路径是可开关的：如果模型训练时没有用 Engram 配套的压缩词表，推理端就退化为恒等映射。

值得注意的是 cfg.enable_tokenizer_compression 这个开关也影响着哈希侧的词表大小：

1
vocab_size = (
2
    len(CompressedVocabulary(cfg.tokenizer_name_or_path))
3
    if cfg.enable_tokenizer_compression else int(cfg.tokenizer_vocab_size)
4
)

也就是说，压缩开关必须与训练时的配置完全一致。改变它会同时改变 lookup_table 和哈希的取模上界，导致所有表行错位。这类「训练-推理一致性」问题在推理框架里很常见，但 Engram 把它推到了极致：不是权重对不上，而是地址对不上，而且不会触发任何形状检查，只会静默地取回错误的向量。

三层参数：阶数、头、层

Engram 的哈希不是「一句话一个哈希」，而是按 N-gram 阶数与哈希头展开的多组哈希。EngramLayout.from_cfg 把这些维度固定下来：

1
@dataclass(frozen=True)
2
class EngramLayout:
3
    max_ngram_size: int
4
    layer_ids: list[int]
5
    layer_vocab_sizes: dict[int, list[list[int]]]
6
    ngram_embed_dim: int
7

8
    @classmethod
9
    def from_cfg(cls, cfg: AttrDict) -> EngramLayout:
10
        max_ngram_size = int(cfg.engram_max_ngram_size)
11
        assert max_ngram_size >= 2 and cfg.engram_layer_ids
12
        assert len(cfg.vocab_size) == len(cfg.n_embed_per_ngram) \
13
            == len(cfg.n_head_per_ngram) == max_ngram_size - 1
14
        num_heads = int(cfg.n_head_per_ngram[0])
15
        assert num_heads > 0 and all(h == num_heads for h in cfg.n_head_per_ngram)
16
        assert all(dim % num_heads == 0 for dim in cfg.n_embed_per_ngram)
17
        head_dim = cfg.n_embed_per_ngram[0] // num_heads
18
        assert all(dim // num_heads == head_dim for dim in cfg.n_embed_per_ngram)

这段断言把配置的自由度收得很紧：

• max_ngram_size = N：最多看 N 个 token（当前 token 加 N−1 个回看）。
• vocab_size、n_embed_per_ngram、n_head_per_ngram 的长度都是 N - 1，对应阶数 2 到 N 的 N−1 个 N-gram 表。1-gram（单 token）不做表，因为它等价于普通词嵌入。
• 每个阶数的头数必须相同，每个阶数的嵌入维度必须能被头数整除，并且所有阶数的 head_dim 必须一致。原因很实际：每个 token 的 N-gram 记忆向量是把各阶、各头的结果拼起来，下游 fused_kv_projection 的输入维度是固定的 (N-1) × K × head_dim。如果某个阶数的 head_dim 不同，拼接就需要补零或额外切分，得不偿失。

素数模数：为什么每张子表都是不同的素数

from_cfg 的后半段是分配表大小：

1
seen: set[int] = set()
2
layer_vocab_sizes: dict[int, list[list[int]]] = {}
3
for layer_id in cfg.engram_layer_ids:
4
    per_ngram = []
5
    for vocab_size in cfg.vocab_size:
6
        start = int(vocab_size) - 1
7
        sizes = []
8
        for _ in range(num_heads):
9
            prime = int(nextprime(start))
10
            while prime in seen:
11
                prime = int(nextprime(prime))
12
            seen.add(prime)
13
            sizes.append(prime)
14
            start = prime
15
        per_ngram.append(sizes)
16
    layer_vocab_sizes[layer_id] = per_ngram

三个设计点：

1. 取素数。合数模数会让乘性哈希的低位模式聚集到有限的余数类上。例如模数含因子 2，奇数乘子与偶数乘子会把输入分到不同的奇偶类；模数含小质因子时，低位相同的输入容易撞在一起。素数模数没有非平凡因子，乘法混合后的余数分布更均匀。
2. 所有层、所有阶数、所有头的素数互不相同。这一点比「每张表各自取素数」更强。因为最终所有子表被拼成一张扁平的大表，用 offsets 记录每张子表的起始行；如果两张子表的素数相同，偏移就只相当于给哈希结果加了一个常数，两个头的碰撞模式会高度相关。让每个头的模数都不同，等于让每个头的哈希函数在整数环上「错开」，碰撞事件互相独立。
3. 素数从 vocab_size - 1 开始向上取。这个上界决定了表的大小与词表规模同量级，不会出现「词表十万、表四亿」的荒谬配置；真正的「四亿行」来自多个阶数、多个头、两个层以及训练时选的 vocab_size 共同累加。

seen 集合是全局的，跨层也生效。这带来一个副作用：后一层的素数会比前一层更大一些，表行数也略多。V4.1-Flash 的两层表分别约 3.840 亿与 3.840 亿行，差距就来自这里。

奇数乘子：把单个 token 扩散到整个字长

NgramHashMapping 为每一层生成乘子：

1
bound = max(1, (np.iinfo(np.int64).max // vocab_size) // 2)
2
multipliers = []
3
for layer_id in self.layer_ids:
4
    rng = np.random.default_rng(int(seed + 10007 * int(layer_id)))
5
    multipliers.append(
6
        rng.integers(0, bound, size=(layout.max_ngram_size,), dtype=np.int64) * 2 + 1
7
    )
8
self.ngram_multipliers = torch.tensor(np.stack(multipliers), dtype=torch.int64, device=device)

• 乘子为奇数（r * 2 + 1）。奇数在模 2^k 下可逆，乘法的低位不会因为因子 2 而塌缩；配合奇素数模数，乘性混合能把输入 id 的差异扩散到结果的高位。
• 每层不同的随机种子（seed + 10007 * layer_id）。如果两层用同一组乘子，同一对 N-gram 在两层里会以相同方式碰撞，错误会在两层重复叠加；不同乘子让两层的碰撞模式独立，模型看到的至少有一层是正确的。
• 上界 bound = i64max / vocab_size / 2。这是为了防止 token * multiplier 在 int64 里溢出。token 最大是词表大小，乘子最大是 bound，乘积不超过 i64max/2；后面的异或与累加也不会越过 int64。这个界不是随便写的，它把「哈希可用」与「算术溢出」之间的安全边界固定下来。
• 每个阶数一个乘子。N-gram 的每个位置有自己的权重，位置信息不会被平均掉。否则「AB」与「BA」会得到相同的混合结果。

engram_metadata.py 只负责生成参数，真正的混合公式在内核 tile_kernels.engram.engram_hash 里。论文给出的形式是乘法-XOR：

其中：

• $n$ 是 N-gram 阶数，取值 $2 \ldots N$；
• $k$ 是哈希头编号，取值 $0 \ldots K-1$；
• $\text{token}_{q-r}$ 是第 $q$ 个 query 往回看 $r$ 步的压缩 token id；
• $m_r$ 是该层第 $r$ 个位置的奇数乘子；
• $p_{n,k}$ 是阶数 $n$、头 $k$ 的素数表大小；
• $o_{n,k}$ 是这张子表在扁平大表中的起始行偏移，由 tile_kernels.torch.engram.make_offsets 从 ngram_vocab_sizes 算出；
• 异或 $\oplus$ 是按位异或，不是加法。

异或代替加法的原因是：加法有进位，低位的变化会串扰到高位，混合结果对输入的低位过于敏感；异或是不带进位的按位混合，每个比特独立扩散，配合乘法的「乘法把差异铺满整个字长」，整体分布更均匀。

一次哈希算出多少个索引

每个 token 在每个 Engram 层要查 (N-1) × K 个表项：阶数 2 到 N 各 K 个头。以 $N=4$、$K=8$ 为例，每个 token 每层 24 个索引，两层就是 48 个。索引数组的形状是：

1
hashes: [num_layers, num_tokens, num_entries_per_token]

这正是 EngramBuffer.fetch(indices) 期望的形状。fetch 的文档写得很明确：indices 是 [num_layers, num_tokens, num_entries_per_token]，一个 token 的所有表项一次性发出，按层组织，便于内核按层遍历。

索引是全局行号，不是「第几个 rank 的第几行」。全局行号的范围是 [0, num_embeddings)，其中 num_embeddings = sum(所有阶数、所有头的素数)。远程取数内核拿到全局行号后，先减去本 peer 的起始全局行号，再除以每 rank 行数得到 rank，最后取模得到行内偏移。把「路由」放在内核里做，而不是在 Python 里按 rank 分组，避免了一次大规模排序与多次小批量调用。

图像位置的处理

engram_hash 的最后一个参数是 images_seq_mask。图像 token 在压缩空间是 -1，它既没有文本语义，也不应该参与 N-gram。内核会根据 mask 把这些位置强制映射到 pad 行（或跳过），避免为视觉 token 发起无意义的远程读取。这里不展开具体分支，只需记住：图像 token 不产生有效的记忆检索，这与后文 n-gram 窗口在图像边界处截断是同一套语义的两个侧面。

按行切：行数与 ENP 的关系

NgramEmbeddingParameters 负责把一层的表按 ENP（Engram 并行）切开：

1
head_vocab_sizes = tuple(size for sizes in layout.layer_vocab_sizes[layer_id] for size in sizes)
2
embedding_dim = layout.ngram_embed_dim
3
alignment = math.lcm(256, engram_parallel_size)
4
num_embeddings = sum(head_vocab_sizes)
5
padded_rows = (num_embeddings + alignment - 1) // alignment * alignment
6
num_per_part = padded_rows // engram_parallel_size
7
row_start = engram_parallel_rank * num_per_part
8
assert row_start < num_embeddings
9
self.embedding = nn.Parameter(torch.empty((num_per_part, embedding_dim), dtype=torch.bfloat16))

几个关键量：

• head_vocab_sizes 把阶数与头两维拉平成一维。一张层内表实际上由 (N-1) × K 张子表首尾相接组成，行号 global_idx 在这张扁平的层内表上编号。
• alignment = lcm(256, ENP)。行数先向上对齐到 256 与该 ENP 大小最小公倍数的整数倍，再平均分配。为什么要对齐 256？因为 FP8 的 scale factor 是按 32 列或 128 列一块打包的（取决于平台），如果每层的表行数不是块大小的整数倍，边界的 FP8 块就会跨层或跨 rank 被切开。对齐到 256 是一个安全的下界，保证任何常见的块大小都能整除。
• num_per_part 必须使每个 rank 的行数相同。哈希结果落在哪个 rank 是不确定的，如果某个 rank 的行数不同，内核里「global_idx / num_per_part 得到 rank」这一步就不成立。所以源码选择了「先补行、再均分」，补出来的行是永远不会被哈希命中的 padding 行。
• row_start 被挂在参数上（embedding.row_start = row_start），供加载时做覆盖校验与按 rank 物化。上一篇（第十四篇，权重装载）已经讲过这个机制，这里只需要知道：每个 rank 只加载自己那一段行，不加载整张表。

NgramEmbedding 构造时会断言 weight.shape[0] == num_local_embeddings，即「本地权重行数必须等于按该 ENP 大小切出来的行数」。这是防止用 A 并行度切出来的权重去喂 B 并行度的运行时，属于典型的检查-使用一致性问题。

一张表有多大

以 V4.1-Flash 的配置量级估算：两层，每层约 3.84 亿行，每行 256 维。若以 BF16 存储，单层约

两层合计接近 400 GB。这个量级是本节所有工程决策的根源：

• 任何单个 GPU 都装不下；
• 即使切成 64 份，每个 rank 也有约 3 GB；
• 放到 CPU 内存里，则要求每个进程的可用内存足够容纳自己的分片，并支持 RDMA 注册。

DeepInfer 的选择是：默认存 CPU 内存，通过 RDMA 读。userbuffer.allocate_engram 里有一行：

1
use_cpu_rdma_storage = not envs.DEEPINFER_ENGRAM_FORCE_GPU_STORAGE

并有一句注释说明了这个选择的意义：把表放在 GPU 显存里显然不现实。留一个强制 GPU 存储的环境变量，是为了小模型或大显存场景下做对照实验，而不是生产默认。

FP8 表：数据与 scale 分开存

EngramEmbedding 有一个约束 fp8_proj or not fp8_embed：FP8 嵌入表必须搭配 FP8 投影。NgramEmbedding 里对应的部分是：

1
self.fp8_embed = cfg.engram.fp8_embed
2
self.weight = source_embedding.embedding.data
3
self.sf = getattr(source_embedding.embedding, "sf", None)

FP8 行有 256 维，按 32 维一组切 block，共 8 个 block，每个 block 一个 scale。scale 采用 ue8m0 格式（8 位、无符号指数、无尾数，即「2 的幂」），这是 MXFP8 的 scale 约定：数值 = 2^(exp - 127)，只占 1 字节。

scale 不是跟着行一起散落在同一个张量里，而是单独一张 sf 表，并被打包成整数传输：

1
def _pack_ue8m0_to_int(x: torch.Tensor) -> torch.Tensor:
2
    """Pack the ue8m0 scales represented in fp32 into int32 (CUDA) / int16 (ASC), 4/2 per group."""
3
    assert x.dtype == torch.float32 and x.size(-1) % 4 == 0
4
    x_int32 = x.view(torch.int32)
5
    assert (x_int32 >= 0).all() and (x_int32 & 0x7FFFFF == 0).all()
6
    return torch.floor_divide(x_int32, 1 << 23).to(torch.uint8).view(PLATFORM.fp8_sf_type)

这里有两个值得讲透的细节：

1. 为什么用 fp32 表示再打包。ue8m0 的数值只有 256 种，但很多中间工具会用 fp32 保存一个精确的 2 的幂，以便和别的 scale 做统一处理。打包时断言「指数位以内、尾数全零」，然后用 >> 23 取出指数，得到 8 位的 ue8m0。
2. 4/2 个一组打包。CUDA 平台把 4 个 ue8m0 打进一个 int32，Ascend 平台打进 int16。分组数必须满足 groups % itemsize == 0，这也是 EngramRuntime.allocate 里断言 n_embed_per_head % 32 == 0 的原因：256 / 32 = 8 个 block，8 能被 4 整除，也能被 2 整除。

打包的好处是传输与对齐：远程取数时，一次 RDMA 读拿到的不是 8 个孤立的字节，而是 2 个 int32（CUDA）或 4 个 int16（Ascend），对齐到 4/2 字节的自然边界。内核里把整数值写到接收区后，GEMM 一侧把它 view 回 ue8m0 数组即可。

gather_sf 在初始化时把所有 rank 的 scale 拼起来：

1
local_sf = _pack_ue8m0_to_int(self.sf)
2
sf = local_sf.new_empty(self.engram_group.size() * self.num_local_embeddings, local_sf.size(1))
3
dist.all_gather_into_tensor(sf, local_sf, group=self.engram_group)
4
return sf

为什么 scale 要在初始化时 all-gather 一份全量？因为数据面走的是 RDMA，而 scale 走的是 NVLink/LSA（load-store 访问）：DeepEP 的取数内核里，FP8 的 scale 用 gin.get_sym_ptr<ncclTeamTagLsa> 直接读对等 rank 的对称内存，不需要 RDMA 往返。这意味着每层的 scale 表（行数 × 8 字节）在初始化时被复制到所有参与 ENP 的 rank 上，容量是 行数 × 8 B。按四亿行算，约 3 GB；如果 scale 也放在 CPU 内存里，就无法用 LSA 直接读，只能再走一次 RDMA，而 scale 的访问频次与数据行相同，这会把小包 RDMA 的数量翻倍。

一个容易忽略的结论：FP8 省的是数据流量，并不省 scale 的存储。当 ENP 很大时，scale 的全量复制会成为显存约束的一部分，配置时要把这笔账算进去。

显存侧要预留多少

EngramBuffer.get_storage_size_hint 返回两个数：主 GPU 缓冲大小与每 rank 的 RDMA 存储大小。文档说明 GPU 侧的大小「覆盖取数接收区加上本 rank 那份 FP8 scale 表」。也就是说，运行时 GPU 上占用的是：

第一项中的 $T$ 是该引擎每步最多处理的 token 数（decode 用 max_num_decoding_tokens，prefill 用 max_num_prefilling_tokens），乘 2 是因为 fp8_embed 为真时每行还对应一份 scale 字节。

这里可以看到一个重要的取舍：接收区是每步预算，而不是每层独立。所有层、所有 token、所有表项的结果被连续写进同一块接收区，每层的 hook 只是在这块大缓冲区里指向自己的那一小段。这样做的好处是一次分配、零碎片，并且能够和 CUDA Graph 配合——接收区的地址必须稳定，图中才能记下正确的指针。

行对齐与 ENP 的另一个约束

EngramLayout 与 NgramEmbeddingParameters 都有一个隐含前提：所有 rank 看到的 layout 完全一致。素数在配置顺序上依次分配，如果某个 rank 的 engram_layer_ids 顺序不同，seen 集合的消费顺序就会不同，后一层的素数会变。DeepInfer 的 builder 从同一份 config 构造 layout，不依赖 rank 号，因此天然一致。但自定义部署时如果有人手工改过某个 rank 的 config，就会出现「同一 N-gram 在不同 rank 映射到不同素数」的问题，而且只表现为输出漂移，不会报错。

为什么需要一张「按需取回」的缓冲

先做一个对比。MoE 的 dispatch/combine 是「把 token 搬到专家所在的 rank」。Engram 的语义不同：每个 token 都要读若干个固定行，行落在哪个 rank 由哈希决定，读出来的结果是该 token 的「记忆向量」。

• 如果沿用 dispatch 语义，就要把 token 按目标 rank 分组、发出去、再收回来，一次 N-gram 查询会变成多轮 all-to-all，而且接收端要做「按 token 重新拼接」。
• 如果采用 gather 语义，目标 rank 只需要把被命中的行「点名」读到发起端的缓冲里，发起端按原顺序拿到结果，不需要额外重排。

Engram 采用的是后者。发起端是 GPU 上的一个 kernel，它持有索引数组与一块接收缓冲；对端不需要任何计算，只需要允许远端直接读取自己的一块注册内存。这正是 RDMA 的单边读（one-sided read）语义，也是 DeepInfer 复用 DeepEP 的原因——DeepEP 已经把 NCCL 的对称内存、QP 管理与 GPUDirect 这条链路封装好了。

EngramBuffer：一个专门为查表封装的缓冲

deep_ep.buffers.engram 里的 EngramBuffer 与 MoE 用的 Buffer 共享 BufferBase，但配置完全不同：

1
class EngramBuffer(BufferBase):
2
    """GPU fetch buffer with separate CPU or GPU storage shards registered for RDMA."""
3

4
    def __init__(self, group, num_gpu_bytes, num_rdma_storage_bytes,
5
                 use_cpu_rdma_storage, num_allocated_qps, qp_depth,
6
                 restrict_rd_atomic=False, allow_hybrid_mode=True, ...):
7
        ...
8
        num_registered_bytes = (num_gpu_bytes + num_rdma_storage_bytes * num_max_local_ranks
9
                                + (1 << 32))
10
        num_total_gpu_bytes = torch.cuda.get_device_properties('cuda').total_memory
11
        if num_registered_bytes > num_total_gpu_bytes:
12
            os.environ['NCCL_WIN_STRIDE'] = str(align(num_registered_bytes, 1 << 32))
13
        ...
14
        self.runtime = _C.EngramBuffer(
15
            self.rank_idx, self.num_ranks, self.nccl_comm_handle.get(), shared_comm,
16
            num_gpu_bytes, num_rdma_storage_bytes, use_cpu_rdma_storage,
17
            allow_hybrid_mode, sl_idx, num_allocated_qps, qp_depth, ...)

逐行看几个关键设计：

1. 两块内存分开放：num_gpu_bytes 是 GPU 上的接收区与 scale 副本；num_rdma_storage_bytes 是每 rank 存表分片的区域，可能在 CPU（默认）也可能在 GPU。两者都在 NCCL 对称内存窗口里注册，但用途不同。
2. NCCL_WIN_STRIDE 的调整：当注册的总量超过本机显存容量时，说明窗口不可能以「每 rank 一份完整副本」的 stride 映射；NCCL_WIN_STRIDE 改为按总量对齐到 4 GB 的方式设置。这是 CPU 存储路径能成立的前提：window 的虚拟地址空间覆盖所有 rank 的 CPU 分片，但物理内存分散在各进程里。
3. allow_hybrid_mode：每个节点只需注册本节点的存储分片，节点间用 RDMA，节点内用 NVLink/LSA。EP_NUM_MAX_LOCAL_RANKS 默认 16，num_registered_bytes 里乘 num_max_local_ranks 就是因为一个节点上要注册多份。
4. create_shared_handle：CPU 存储时，把 (pid, fd) 通过 all_gather_object 发到组内每个 rank，让同节点的其它进程可以 mmap 同一份 CPU 内存。这样同一节点内一个 rank 的 CPU 分片对所有本地 rank 可见，不需要走网络。
5. 强制新的 NCCL communicator（force_new_comm=True），并且只注册 GDAKI 后端（NCCL_GIN_TYPE=3）。MoE 的通信与 Engram 的查表在流量特征上完全不同——前者是突发的大包，后者是持续的 128 字节小包——分开 communicator 可以让 QP 与 SL（服务等级）独立配置。
6. restrict_rd_atomic 与 NCCL_GIN_GDAKI_MAX_QP_RD_ATOMIC：限制每 QP 的读原子数量，避免网卡内的读请求队列成为瓶颈。

QP 的数量是怎么定的

RDMA 的吞吐不仅取决于带宽，还取决于「有多少请求同时在途」。DeepEP 用 get_theoretical_config 按理论模型估算需要多少 QP：

1
@staticmethod
2
@functools.lru_cache(maxsize=None)
3
def get_theoretical_config(group, num_layers, num_max_tokens,
4
                           num_entries_per_token, num_max_read_requests_per_qp=16):
5
    nic_read_mpps = 90 if check_fast_rdma_atomic_support() else 45
6
    return _C.EngramBuffer.get_theoretical_config(
7
        comm.get_nccl_comm_handle(group).get(),
8
        num_layers, num_max_tokens, num_entries_per_token,
9
        nic_read_mpps, num_max_read_requests_per_qp)

这个函数的三个输入正好对应表的访问量：层数、最大 token 数、每 token 的表项数；输出 (num_qps, qp_depth, restrict_rd_atomic)。nic_read_mpps 是「网卡每秒能发多少个读请求」的估算，硬件支持快速 RDMA 原子操作时为 90 Mpps，否则为 45 Mpps。

可以用一个粗算理解为什么 QP 数量不可忽略。以两步 decode 的典型规模为例：64 个序列、每步 6 个 query（nextn=5）、每 token 每层 24 个表项、两层，则每次 fetch 的读请求数为

每个读请求读 256 字节，总数据量约 4.7 MB。如果这些读请求按顺序从单个 QP 发出，即使每请求 0.1 µs 的发射开销也要接近 1.8 ms，远超一次 decode 的计算时间。多 QP 并行发射、每个 QP 维持多个在途请求，才能把发射时间压到与计算重叠的范围内。这是「小包为主」的数据面与 MoE「大包为主」的数据面最大的区别之一，也是 Engram 必须有自己的 QP 配置的原因。

数据面：engram_fetch_impl 一次取数做了什么

真正的读逻辑在 DeepEP 的 deep_ep/include/deep_ep/impls/engram/engram_fetch.cuh。这个 kernel 用模板参数把几乎所有决定都编译期固定下来：

1
template <int kNumSMs, int kNumQPs,
2
          int kNumSFWarpsPerSM, int kNumIssueWarpsPerSM,
3
          typename engram_layout_t,
4
          int kNumSFPacks,
5
          int kNumMaxTokens,
6
          int kNumRDMAPeers,
7
          int kNumRanksPerRDMAPeer,
8
          int kFlushDepth,
9
          int64_t kNumTimeoutCycles,
10
          ...>
11
__global__ void __launch_bounds__(kNumThreads, 1)
12
engram_fetch_impl(
13
    const ncclDevComm_t nccl_dev_comm,
14
    const ncclWindow_t nccl_window,
15
    const ncclWindow_t* storage_windows,
16
    void* fetched,
17
    void* workspace_ptr,
18
    int* indices,
19
    ncclGinRequest_t* last_gin_requests,
20
    const sf_pack_t* sf_shards,
21
    sf_pack_t* fetched_sf, ...
22
    const int num_tokens);

参数里的每一项都在说明这个 kernel 的形态：

• kNumRDMAPeers、kNumRanksPerRDMAPeer：把对端按「RDMA peer」分组，每个 peer 下挂若干个 rank。一个 NVLink 域内的多个 rank 共享一个 RDMA peer 的网卡通道。
• kNumQPs：同一 peer 下的 QP 数量，issue warp 按 (peer, QP) 二维展开。
• storage_windows：每个对端 rank 的存储窗口指针，索引到某个 rank 后直接用它做 GIN 的源地址。
• fetched：本端的接收缓冲；fetched_sf：scale 的接收缓冲。
• workspace_ptr：一块用于发射计数与同步的工作区。
• last_gin_requests：每个 (层, peer, QP) 一个「最后一次请求写信号」的位置，用于 flush_async 后等待网卡真正的完成信号。

kernel 按 warp 分成两种角色。

issue warp 负责发出读请求。它的核心是三层循环：外层按层，中层按 token tile（warp 32 个 token 一组），内层按表项展开：

1
if (warp_idx < kNumIssueWarpsPerSM) {
2
    constexpr int kNumPeerContexts = kNumRDMAPeers * kNumQPs;
3
    const auto issue_warp_idx = sm_idx * kNumIssueWarpsPerSM + warp_idx;
4
    const auto num_token_tiles = math::ceil_div(num_tokens, kWarpSize);
5
    const auto issue_layout = EngramIssueWarpLayout<...>(issue_warp_idx);
6
    ...
7
    for (int layer_idx = 0; layer_idx < kNumLayers; ++ layer_idx) {
8
        ...
9
        for (int token_tile_idx = issue_layout.peer_warp_idx;
10
             token_tile_idx < num_token_tiles;
11
             token_tile_idx += issue_layout.num_peer_warps) {
12
            const auto first_entry_idx = token_tile_idx * kWarpSize * kNumEntriesPerToken;
13
            for (int round_idx = 0; round_idx < kNumEntriesPerToken; ++ round_idx) {
14
                const auto entry_idx = first_entry_idx + round_idx * kWarpSize + lane_idx;
15
                const auto global_idx = entry_idx < num_entries
16
                    ? __ldg(indices + layer_idx * num_entries + entry_idx) : -1;
17
                const auto belongs_to_peer =
18
                    global_idx >= peer_global_idx_begin && global_idx < peer_global_idx_end;
19
                const auto peer_mask = __ballot_sync(0xffffffffu, belongs_to_peer);
20
                ...

几个值得讲的细节：

1. peer 的划分。每个 lane 读一个索引，判断它是否落在本 warp 负责的 peer 区间 [peer_global_idx_begin, peer_global_idx_end) 内。这个区间是「若干 rank × 每 rank 行数」拼起来的一段连续全局行号，所以判断只需要两次比较。
2. __ballot_sync + __popc 做 warp 内压缩。属于本 peer 的 lane 把请求紧凑地写进一个环形 pending 缓冲，其它 lane 不占位。这样一批请求可以按 warp 宽度整批发出，而不是逐条发。pending 缓冲是两个 32 槽的环形区，交替使用，保证发批与收集同时进行而不互相覆盖。
3. 门铃（doorbell）聚合。issue_batch 里有一个阈值判断：

1
const auto doorbell_threshold = kFlushDepth >= num_active_peer_context_warps
2
    ? kFlushDepth / num_active_peer_context_warps : 1;
3
const auto rings_doorbell = is_last_batch
4
    or num_issued_requests % doorbell_threshold + batch_size >= doorbell_threshold;
5
const auto options = rings_doorbell ? ncclGinOptFlagsDefault : ncclGinOptFlagsAggregateRequests;

不是每个读都敲一次门铃。多条请求可以先以 AggregateRequests 选项提交，攒到阈值再发一次 Default 让它真正落到网卡上。门铃是控制面的开销，对 256 字节的小包来说，每条都敲会显著降低有效吞吐。

4. 每条请求的寻址：

1
const auto peer_local_idx = global_idx - peer_global_idx_begin;
2
const auto intra_peer_rank_idx = peer_local_idx / num_entries_per_rank;
3
const auto local_entry_idx = peer_local_idx % num_entries_per_rank;
4
const auto src_byte_offset = engram_layout_t::get_storage_byte_offset(layer_idx, local_entry_idx);
5
const auto dst_byte_offset = engram_layout_t::get_recv_byte_offset(layer_idx, entry_idx, kNumMaxTokens);
6
gin.get<ncclTeamTagWorld, ncclCoopThread, ncclGin_SegmentMixed>(
7
    storage_windows[intra_peer_rank_idx], src_byte_offset,
8
    nccl_window, gin.get_sym_offset(math::advance_ptr(fetched, dst_byte_offset)),
9
    engram_layout_t::kNumHiddenBytes, rdma_peer_rank, options);

源地址是「对端 rank 号 × 每 rank 行数」的整除取模，目的地址是「(层, token, 表项)」三者决定的固定槽位。两次地址计算都是 engram_layout_t 的编译期函数，算的是字节偏移而不是行号，内核里没有分支、没有除法以外的复杂运算。

注意 entry_idx 是扁平的 token×表项序号，不是行号：同一 token 的 24 个表项在接收区里连续排列，dst_byte_offset 由 (layer_idx, entry_idx, kNumMaxTokens) 决定。这意味着接收区的布局是 [层][token][表项][256 字节]，第 1 层的 hook 拿到的正是这块连续内存里属于第 1 层的那一段。这也是为什么 Engram.forward 里可以把 embed_hook() 的输出直接 unflatten 成 [total_q, 表项数, 256] 再进投影——布局与语义是一一对应的。

5. flush 与最后一拍的窗口。所有请求的 ncclGinRequest_t 由同 QP 上最后一个完成的 warp 负责 flush_async：

1
if (num_active_peer_context_warps == 1) {
2
    gin.flush_async<ncclTeamTagWorld, ncclCoopThread>(rdma_peer_rank, request_ptr);
3
} else {
4
    cuda::atomic_ref<uint32_t, cuda::thread_scope_device> num_arrivals(*issue_counter);
5
    const auto arrival_idx = num_arrivals.fetch_add(1, cuda::memory_order_release);
6
    if (arrival_idx + 1 == num_active_peer_context_warps) {
7
        ptx::fence_acquire_gpu();
8
        gin.flush_async<ncclTeamTagWorld, ncclCoopThread>(rdma_peer_rank, request_ptr);
9
        ptx::st_relaxed_gpu(issue_counter, 0u);
10
    } else if (layer_idx + 1 < kNumLayers) { ... 超时等待 ... }
11
}

flush_async 是告诉网卡「已提交的这批请求现在发出去」。一个 QP 上有多个 issue warp 时，用 device 作用域的原子计数确保所有 warp 都提交之后才敲最后一次门铃；下一个 warp 在进入下一层之前等待计数归零。这里的等待是设备内同步，不涉及主机，不会打断 host 端已经发出的后续 kernel。超时分支会打印 DeepEP Engram issue barrier timeout，这是排查「某一层总比别的层慢」时第一个要看的信息。

scale warp 负责搬 FP8 的 scale。它走的是 ncclTeamTagLsa 的内存直接访问，不是 RDMA：

1
const auto* sf_shard = gin.get_sym_ptr<ncclTeamTagLsa>(const_cast<sf_pack_t*>(sf_shards), sf_shard_idx);
2
const auto* src = reinterpret_cast<const uint64_t*>(sf_shard + sf_row_idx * kNumSFPacks);
3
uint64_t values[kNumVecs];
4
#pragma unroll
5
for (int i = 0; i < kNumVecs; ++ i)
6
    values[i] = __ldg(src + i);
7
auto* dst = fetched_sf + layer_idx * sf_layer_stride +
8
    (entry_idx / kNumEntriesPerToken) * sf_token_stride +
9
    (entry_idx % kNumEntriesPerToken) * kNumSFPacks * sf_hidden_stride;

每个 entry 的 scale 是 kNumSFPacks 个 ue8m0 字节（256 维、32 维一组 = 8 个字节，CUDA 打包后 2 个 int32）。这里用 uint64_t 一次读 8 字节，再按 sf_hidden_stride 列主序散写到接收区。列主序（column-major）与 TMA 对齐是 GEMM 侧的要求：投影矩阵消费 scale 时按列读，必须保证每个 block 的 scale 连续。

为什么 scale 不走 RDMA？因为它小。每行 8 字节，一次 fetch 的总 scale 量是 token × 24 × 8 ≈ 72 KB（64 序列、6 token 的场景），走 LSA 的 load 比走网卡更便宜，而且 LSA 访问的是所有 rank 的 scale 副本（初始化时 all-gather 过），不需要按 rank 定位——这省掉了「找对端 rank 的窗口指针」这一步。

回看：fetch 返回给 Python 的是什么

EngramRuntime.fetch 调用 self._buffer.fetch(hashes, use_tma_aligned_col_major_sf=True)，返回 {layer_id: hook}。hook 是引擎侧的一个可调用对象，它内部会做「等数据到齐」的同步——在 CUDA 语义里，读结果的 kernel 与发出请求的 kernel 在同一个 stream 上排队，所以「等」不是 CPU 阻塞，而是 stream 顺序保证。

这一点对理解整个设计至关重要：Python 在模型前向开始时调用 fetch，此时只是发动了读请求；真正读数据的 kernel 在第 1 层才被发射。中间夹着的第 0 层 kernel 与第 1 层前半段 kernel 在 GPU 上执行，网卡读数据与 GPU 计算并行。这就是「把传输藏在计算里」的具体实现。

问题：n-gram 窗口从哪几个 token 取

在非投机、非 ZOS 的朴素实现里，第 $q$ 个 query 的 n-gram 窗口就是「当前位置往回数 $r$ 个 token」，$r = 0 \ldots N-1$。但在 DeepInfer 的服务模式下，三个机制叠加在一起，让这件事变得不平凡：

• 投机解码：一步输入 nextn + 1 个 token，其中 1 个是上一拍采样的 token，nextn 个是草稿 token。草稿 token 是否被接受要等验证之后才知道，但 n-gram 必须在验证前算出来，否则查表来不及藏在计算里。
• ZOS（零开销调度）：上一拍的采样结果与接受数在下一拍才处理，因此「上一拍实际接受了几个草稿」在这一步开始时是知道的（通过 prev_num_accepted_tokens），但调度器原本的 seqlens 是乐观估计。
• 多序列批处理：一个 batch 里有序列是「本步新进来的」（fresh），也有序列是「上一步已经在跑的」（prev），两者的 token 来源不同。

gather_engram_tokens_zos 的接口

DeepInfer 用 tilelang_ops/gather_engram_tokens_zos.py 里的 gather_engram_tokens_zos 一次性解决这三件事。接口签名很长，但每一项都有明确的语义：

1
def gather_engram_tokens_zos(
2
    win_compressed: torch.Tensor,        # [n_seq, window_size] int32, 压缩后的历史（CPU 构造）
3
    curr_fresh: torch.Tensor,            # [n_seq, max_seqlen_q] int32, fresh 序列本步压缩 token（CPU 构造）
4
    prev_sampled_token_ids: torch.Tensor, # [prev_n_seq] int32
5
    prev_predict_token_ids: torch.Tensor, # [prev_n_seq, nextn] int32
6
    prev_num_accepted_tokens: torch.Tensor, # [prev_n_seq] int32
7
    remap: torch.Tensor,                 # [n_seq] int32，-1 表示 fresh 序列
8
    lookup_table: torch.Tensor,          # [vocab] int32，GPU 侧压缩映射
9
    seqlens_q: torch.Tensor,             # [n_seq] int32
10
    cu_seqlens_q: torch.Tensor,          # [n_seq] int32
11
    out_ngram_token_ids: torch.Tensor,   # [total_q, max_ngram_size] int32
12
    *, max_seqlen_q: int, max_ngram_size: int,
13
    engram_compressed_pad_token: int, image_sentinel: int,
14
) -> None:

函数 docstring 里的规则是整个第五段的核心，把它逐条拆开：

其中：

• $i$ 是序列下标，$j$ 是该序列内的 query 下标（$0 \ldots \text{seqlens\_q}[i]-1$），$r$ 是回看步数（$0 \ldots N-1$）；
• $q_{base}$ 是该序列在本 batch 扁平 token 中的起始位置（cu_seqlens_q[i-1]，$i=0$ 时为 0）；
• $p_{\text{idx}} = \text{remap}[i]$，$\ge 0$ 表示该序列是上一拍延续下来的，值是它在上一拍 batch 里的位置；
• $\text{nr} = \text{nextn} - \text{prev\_num\_accepted\_tokens}[p_{\text{idx}}]$（fresh 序列取 $\text{nextn}$），表示上一拍未被接受的草稿数；
• $W$ 是窗口长度 window_size。

逐个分支解释：

分支一（fresh，$j \ge r$）。新序列本步的 token 已经在 CPU 侧压缩好，直接按「本步内第 $j-r$ 个」取。当 $j-r=0$ 时就是当前 token 本身。

分支二（prev，$j \ge r$）。延续序列本步的 token 由上一拍的采样与草稿产生：第 0 个是上一拍的真实采样 token（prev_sampled_token_ids），第 $j$ 个是第 $j-1$ 个草稿 token（prev_predict_token_ids）。这些是原始 id，需要用 GPU 侧的 lookup_table 转成压缩 id。注意即使是草稿 token 也照常进入 n-gram——这一步发生在验证之前，草稿一旦被拒绝，对应的若干行 n-gram 本来就是多余的，但代价只是多读了几行表，换来的是把查表提前到与计算重叠的位置。这是一个有意为之的取舍。

分支三（$j < r$）。查询比回看步数更靠前时，窗口跨到「上一拍之前」。这里的关键是 $W-\text{nr}+j-r$（docstring 里写作 window_size - num_reject + j - r）。为什么是 $W-\text{nr}$？

把压缩历史想象成一条虚拟序列，尾部结构是：

1
compressed_seq_data   last_next_n_token      curr tokens
2
[0, L)                [L, L+nextn)           [L + nextn, L + 2*nextn + 1)
3
              |<---- window_size ---->|

win_compressed 从位置 $L$ 开始取 $W = \text{nextn} + N - 1$ 个 token，覆盖「上一拍的 nextn 个草稿」与「本步的前 $N-1$ 个 token」。如果上一拍有 nr 个草稿被拒绝，那么真正有效的历史末尾就往前移了 nr 个位置——窗口的起点应该相应往后挪 nr，即 base = window_size - nr。这样 query $j=0$（上一拍采样出的 token）回看 $r$ 步时，落在 win_compressed[base - r]，正好是真实历史里的倒数第 $r+1$ 个 token。

可以用一个例子校验：nextn = 5，N = 4，W = 8。上一拍 5 个草稿全部被接受（nr = 0），base = 8。本步第 0 个 query 回看 1 步，应该看到「上一个采样 token」，即上一拍最后一个被接受的草稿，在窗口里是位置 8 - 1 = 7，也就是 win_compressed 的最后一位。上一拍只接受了 2 个（nr = 3），base = 5，第 0 个 query 回看 1 步落在位置 4，即本步第一个 token 往前一位的「上一拍第 2 个被接受的草稿」。两种情况下都对。

如果漏掉 $-\text{nr}$，窗口会把已拒绝的草稿也算进去，模型看到的 N-gram 会在被拒绝的位置上错误地拼接。这种错误不会崩，只会让 Engram 在投机解码路径上取到错误表项，进而造成输出静默漂移——这也是为什么这一块必须有 CPU 参考实现 gather_engram_tokens_zos_ref 与参数化单测。

图像哨兵：为什么「一旦回看到图像，该阶及以上全部置 pad」

规则的最后一段是 VL（视觉语言）模型特有的：

图像 token 在 compressed 空间是哨兵 -1；沿 lookback 维 r 一旦回看到 -1，该 r 及更高阶置 engram_compressed_pad_token。

直觉上很好解释：图像 token 是一个占位符，真实内容在 image_seq_mask 标记的位置，不参与文本 N-gram。如果一个 N-gram 窗口里包含图像 token，这个 N-gram 的文本意义就是断裂的——「前文 + 一张图 + 后文」不是一个自然的词语搭配。与其给它一个错误的一阶组合，不如直接让它失效，映射到 pad 行。

但「一旦回看到就全部置 pad」有一个更强的含义：高阶 N-gram 的失效范围比低阶更大。比如第 0 个 query 回看 2 步时遇到了图像，那么 $r=2$ 与 $r=3$ 都失效，$r=0$、$r=1$ 仍然有效。换言之，图像位置附近仍然保留短程 N-gram，只是长程 N-gram 被切断。这与「图像是一个语义边界」的直觉一致，也避免了在图像周围制造整片空窗口。

单元测试告诉我们的约束

gather_engram_tokens_zos.py 自带了一个 CPU 参考实现和参数的矩阵测试：

1
@pytest.mark.parametrize("old_ratio", [0.0, 0.5, 1.0])
2
@pytest.mark.parametrize("nextn,max_ngram_size",
3
    [(0, 3), [0, 4], (1, 3), (1, 4), [5, 3], [5, 4]])
4
def test_gather_engram_tokens_zos(nextn, max_ngram_size, old_ratio):
5
    ...
6
    out_ref = gather_engram_tokens_zos_ref(...)
7
    assert torch.equal(out_gpu.cpu(), out_ref)

测试覆盖了四个维度：是否有投机（nextn 0/1/5）、N 的大小（3/4）、prev 序列占比（0/50/100%）、随机插入图像哨兵。这四维正好对应上面讨论的四个分支与两类边界。从测试的覆盖矩阵可以反推出实现里最容易出错的路径：

• old_ratio = 0.0：全是 fresh 序列，走分支一；
• old_ratio = 1.0：全是 prev 序列，走分支二，并且会触发分支三里「草稿被部分拒绝」的小跨；
• nextn = 0：没有投机，W = N - 1，此时分支三只在查询与回看跨过序列起点时触发；
• nextn = 5：nextn 与 N 的交互最强。

还有一条 benchmark 测试值得注意：test_gather_engram_tokens_zos_bench 的参数里固定了 nextn = 5, max_ngram_size = 4, old_ratio = 0.9，这是典型的 decode 负载——几乎全部是延续序列、满投机深度。n_seq 从 64 到 1024 扫描，测的是 kernel 在 batch 规模变化时的线性度。这个 kernel 的复杂度是 O(total_q × max_ngram_size)，每个 (seq, query) 由一个 warp 的 32 个 lane 并行处理 query 维，因此 batch 增大时主要增加的是 SM 的占用而不是单次延迟。

装载侧：表是怎么进到每个 rank 的行上的

前面讲了「运行时按行切表」，但表在磁盘上是一整张。第十四篇（DeepInfer 源码拆解（十四）：权重装载）讲了普通权重与专家权重的三条加载通道，Engram 表是第四条，且单独实现：

1
# deepinfer/hugging_face/weights.py, _load_engram_tables
2
row_start = int(getattr(param, "row_start", 0))
3
with torch.device("cpu"), safe_open(str(hf_dir / shard), framework="pt", device="cpu") as handle:
4
    weight_view, scale_view = handle.get_slice(weight_key), handle.get_slice(scale_key)
5
    valid_rows = min(param.data.shape[0], int(table_shape[0]) - row_start)
6
    rows_per_chunk = max(1, row_chunk_bytes // int(table_shape[1]))
7
    for start in range(0, valid_rows, rows_per_chunk):
8
        stop = min(start + rows_per_chunk, valid_rows)
9
        source = row_start + start
10
        weight = weight_view[source : source + (stop - start)]
11
        param.data[start:stop].copy_(weight)
12
        sf = _ue8m0_to_fp32(scale_view[source : source + (stop - start)].to(device=param.sf.device))
13
        param.sf.data[start:stop].copy_(sf)
14
    if valid_rows < param.data.shape[0]:
15
        param.data[valid_rows:].zero_()
16
        param.sf.data[valid_rows:].fill_(1)

几个与前面段落呼应的点：

• 只读本地行：source = row_start + start，每个 rank 从 HF 全局表的 row_start 处开始读自己那一段，不读整表。这与第三段「按 rank 切表」严格对齐。
• row_start 来自参数本身，而参数在 NgramEmbeddingParameters 构造时被标注（embedding.row_start = row_start）。也就是说，切分方案在模型构造时就决定了，加载只是执行。
• 分块读取，块大小 256 MB（row_chunk_bytes = 1 << 28）。整表在 safetensors 里是连续存储，get_slice 是按行视图读取，不是一次 torch.load。这么大的表如果一次性读入内存，会把主机内存打满；按行切块让峰值内存限定在「一块 + 运行时缓冲」的量级。
• FP8 字节原样拷贝，scale 无损还原：权重保持 1 字节的 FP8，scale 从 ue8m0 字节还原成 fp32 再拷入 param.sf，之后由 gather_sf 统一打包。之所以不在加载时就打包，是因为打包需要知道平台（CUDA 4 个一组 / Ascend 2 个一组）且需要 all-gather，这些都在初始化阶段做更合适。
• 对齐行补齐：HF 表只有 num_embeddings 行，运行时参数有 padded_rows 行（对齐到 lcm(256, ENP)）。超出部分被清零，scale 填 1。填 1 而不是 0 是有意义的：ue8m0 的 1 表示 2^0，如果填 0，某些 GEMM 实现会把 scale 当成 0 处理，导致整行被「静默清零」并可能触发除零保护。对齐行永远不会被哈希命中，所以内容无关紧要，但数值必须「看起来正常」，以免在调试时被当成异常。
• 同 shard 断言：assert _shard_of(index, scale_hf_name) == shard。表和它的 scale 必须在同一个 safetensors 分片里，否则 safe_open 打开的文件里找不到 scale key。这是一个典型的 HF 转换工具需要保证的约定。

与 ZOS 输入准备的关系

prepare_tensor_input_zos.py 里的 prepare_token_ids_zos 是调用方。它把两条路径分开：

• 非 ZOS 路径（metadata_zos.remap is None）：直接用 CPU 侧准备好的 ngram_token_ids_cpu 非阻塞拷贝到 GPU。此时所有序列都是 fresh，没有跨拍修正的问题。
• ZOS 路径：先调用 gather_input_ids_zos 生成本步的 input_ids，再调用 gather_engram_tokens_zos 生成 ngram_token_ids。两者共享同一组 remap、seqlens_q、cu_seqlens_q，保证 input_ids 与 n-gram 窗口的序列排布一致。

这里有一个时序上的细节：ngram_token_ids 在 GPU 上由 kernel 生成，而 win_compressed、curr_fresh 由 CPU 构造后拷入。CPU 侧构造 win_compressed 时用的是压缩后的历史（compressed_seq_data 与上一拍在途 token），这些数据在调度器侧维护。第二篇讲过 ZOS 如何把 CPU 后处理延后一拍，这里看到它的另一个后果：Engram 的窗口维护必须在 CPU 侧把「已提交的压缩历史 + 上一拍未确认的草稿」拼在一起，并在 GPU 上根据接受数做对齐。这两件事分属 CPU 与 GPU，靠 remap 与 prev_num_accepted_tokens 两个小数组衔接。

取回来的不是最终结果，而是投影的输入

EngramBuffer.fetch 的 hook 返回的是什么形状？从接收区布局可以推出来：每层、每 token、每个表项各 256 字节，所以 hook 返回 [total_q, (N-1)×K, 256]。这是一堆原始记忆向量，还差两步才能进入主干：

1. 拼接所有阶数、所有头的向量，经过一个线性投影，得到与主干宽度匹配的表示；
2. 用门控把这个表示融合进主干的多路残差。

Engram.__init__ 把这两步的权重准备好：

1
self.fused_kv_projection = ReplicatedLinear(
2
    engram_module.fused_kv_projection.weight.shape[1],
3
    engram_module.n_embed * (engram_module.hc_mult + 1),
4
    weight=engram_module.fused_kv_projection.weight,
5
)
6

7
# weight_hidden / weight_embed 在推理期恒定，提前融合一次
8
self.weight_fused = tile_kernels.engram.fused_weight(
9
    engram_module.weight_hidden.data,
10
    engram_module.weight_embed.data,
11
)

ReplicatedLinear 就是第十五篇讲过的运行时线性层（DeepInfer 源码拆解（十五）：Linear、RMSNorm 与 RoPE）。这里只关心它的形状语义：

• 输入维度 = sum(n_embed_per_ngram) = (N-1) × K × head_dim，正好等于 hook 返回的扁平表项维度；
• 输出维度 = hidden × (hc_mult + 1)，其中 hc_mult 是 mHC 主干的路径数。

为什么输出是 hc_mult + 1 条路径而不是 1 条？因为 Engram 的输出要写进 mHC 的宽残差主干（关于 mHC 的多路残差结构，可参考站内 mHC 完全拆解）。主干有 hc_mult 条并行路径，记忆向量需要为每条路径提供一份「候选更新」，再加上一份用于门控计算本身的表示，所以投影出 hc_mult + 1 组。

门控：把记忆写进每一条路径

Engram.forward 只有十几行：

1
def forward(self, hidden_states, embed_hook, images_seq_mask=None):
2
    kv = self.fused_kv_projection(embed_hook()).unflatten(
3
        -1, (hidden_states.size(-2) + 1, hidden_states.size(-1))
4
    )
5
    gated, *_ = tile_kernels.engram.engram_gate_fwd(
6
        hidden_states,
7
        kv,
8
        self.weight_fused,
9
        self.eps,
10
        self.clamp_value,
11
        save_for_backward=False,
12
        image_token_mask=images_seq_mask,
13
        inplace=True,
14
    )
15
    return gated

逐项解释：

• hidden_states 的形状是 [total_q, hc_mult, n_embed]。它不是普通 Transformer 的二维 hidden，而是 mHC 主干的三维视图：每个 token 有 hc_mult 条路径。
• unflatten(-1, (hc_mult + 1, n_embed)) 把投影输出的最后一维拆成「路径 × 宽度」。hc_mult 条给主干，最后 1 条留给门控。
• engram_gate_fwd 是一个融合 kernel。它同时做三件事：基于主干与记忆计算门控、对记忆做归一化、把门控后的记忆原地写回 hidden_states。融合的原因很直接：如果拆成三个 kernel，中间结果要落 HBM，total_q × hc_mult × hidden 的读写量在 decode 时并不小；而且三步都需要同一份 hidden_states，融合后只需读一次、写一次。
• self.eps 与 clamp_value = 1e-6。eps 来自模型的 RMSNorm epsilon，用于归一化的数值稳定；clamp_value 则是对门控或归一化分母做下限保护。源码把 clamp_value 固定为 1e-6，不是配置项——说明这是一个经验下界，而不是模型超参。这样做的后果是：当某条路径的激活接近 0 时，门控不会因为除以极小数而爆掉，但也不会精确到「完全关闭」；1e-6 的下限决定了最小可表达的增益。
• image_token_mask。图像 token 的位置不写入记忆。这与第五段 n-gram 在图像边界截断是同一套语义：图像没有文本 N-gram，也就不应该有文本记忆。
• inplace=True。原地写回主干，省掉一份输出缓冲。save_for_backward=False 是因为推理不需要反向图，避免保存中间量。

为什么挂在 hc_post_intervention 上

Engram 层不是一个独立的 decoder layer，而是「在某一层的两个子层之间插一脚」。DecoderLayer.forward 里对应的是这几行：

1
if self.engram is not None:
2
    assert self.hc_enabled and embed_hook is not None
3
    hc_post_intervention = lambda trunk: self.engram(
4
        trunk, embed_hook, getattr(input_metadata, "images_seq_mask", None)
5
    )
6
else:
7
    hc_post_intervention = None
8
...
9
hidden_states, hidden_states_bf16, residual = self.attn_norm_forward(
10
    hidden_states, residual, hc_post_intervention
11
)

s 的位置很讲究：它是 hc_post 的干预点，也就是「上个子层的输出已经写回主干，但当前子层的 pre-mix 与 norm 还没开始」的时刻。放在这里的原因：

• 必须在 hc_post 之后。hc_post 把上个子层的二维输出写回 hc_mult 条路径，只有在这之后，Engram 才能看到完整的主干并同时写入所有路径。如果在子层输出还是二维时就介入，门控只能作用于单条路径，无法表达「不同路径不同更新」。
• 必须在当前子层 pre 之前。当前子层的 pre 会重新混合各路径，Engram 写入的记忆应该参与这次混合；放在 pre 之后就晚了，记忆只能影响一个子层，而不是被后续所有层继承。
• 这也是为什么 assert self.hc_enabled。如果没有 mHC，主干是二维的，hc_post_intervention 这个钩子没有定义位置，Engram 的「多路径写入」语义也无从谈起。

一个被牺牲的融合

HCLayerMixin._hc_fused_post_pre 的正常路径会把「上个子层 post + 当前子层 pre + RMSNorm（+ FP8 cast）」融合进一个 _mega_mhc kernel。但 Engram 层不能走这条路：

1
if (
2
    hc_post_intervention is None
3
    and self.use_mega_mhc
4
    and prev_layer_out.ndim == 2
5
    and prev_lazy_data is not None
6
    and prev_lazy_data.post_layer_mix is not None
7
):
8
    return self._mega_mhc(prev_layer_out, prev_lazy_data, hc_modules, cast_kwargs)
9
...
10
if hc_post_intervention is not None:
11
    next_residual = hc_post_intervention(next_residual)
12
return self._hc_pre(next_residual, hc_modules, prev_pre_layer_mix, cast_kwargs)

注释写得很直白：「no hook can be inserted inside the kernel」。融合 kernel 内部不暴露「post 之后、pre 之前」的中间张量，Engram 只能让这一层退回非融合路径：单独跑 hc_post、调用 Engram、再跑 _hc_pre 与 RMSNorm。代价是两次额外的 kernel 启动，以及中间残差的一次落盘。

这个代价值不值？V4.1-Flash 只有两层 Engram，相比整个模型每步的 kernel 数量，两次启动可以忽略；但如果有人把 Engram 层数加到十几层，融合缺失的代价就会线性增长。源码没有为此做优化（例如把 Engram 的门控也融进 _mega_mhc），原因是门控依赖远程取回的 embedding，而 _mega_mhc 的设计前提是「所有输入都在本地、且形状编译期固定」。远程取数的完成时刻由网卡决定，把等待塞进一个原本纯计算的 mega-kernel 会破坏它的流水线。

投影的 FP8 路径

EngramEmbedding 里的 fused_kv_projection 带一个 use_fp8_gemm=engram_cfg.fp8_proj。当嵌入表也是 FP8 时，取回的行是 FP8 + ue8m0 scale，投影直接吃 FP8 输入，按 block scale 做 GEMM，输出 BF16/FP8 供门控使用。这省掉了一次 FP8→BF16 的反量化：如果表是 FP8 而投影是 BF16，取回的每个 256 维向量都要先转成 BF16，total_q × 24 × 256 个元素的反量化在 decode 路径上会变成实打实的 ALU 与带宽开销。

_pack_ue8m0_to_int 的打包格式与 use_tma_aligned_col_major_sf=True 的接收布局也是为这条路径服务的：GEMM 期望 scale 按列主序、TMA 对齐地排布，接收区就按这个格式写入；如果投影退回 BF16，这些 scale 仍会被取回（因为 fetch 的接口统一），只是不被使用。这是一个「接口统一、按配置裁剪」的设计，读代码时不要误以为所有配置下 scale 都是必需的。

与 mHC、MoE 的对照

把 Engram 的注入方式与 MoE 放在一起看，会发现两者都是「稀疏地改变主干」，但接口完全不同：

这张表的每一行都能在前面的源码里找到对应。最值得记住的是「路由时机」这一行：MoE 的通信无法提前，因为路由依赖运行时激活；Engram 的通信可以提前，因为路由只依赖输入 token。这就是为什么同一套 DeepEP 通信层，在 MoE 里被用作 dispatch/combine，在 Engram 里被用作单边读。

阶段 0：CPU 侧准备（ZOS 修正之后）

ZOS 把上一拍的采样结果与接受数延后了一步。在第一段 GPU 前向之前，Worker 已经完成跨拍修正（第十七篇讲的 _correct_seqlens_qk），并准备好 seqlens_q、cu_seqlens_q、remap、prev_sampled_token_ids、prev_predict_token_ids、prev_num_accepted_tokens。同时，CPU 侧维护的压缩历史 compressed_seq_data 已经把「上一拍已确认的 token」追加进去，并拼出 win_compressed（历史尾部 + 上一拍全部草稿，共 W = nextn + N - 1 个）。

curr_fresh 对全 prev 的 batch 是空的（形状 [n_seq, 6]），只有 fresh 序列才会填。这一步全在 CPU 上做，不涉及 GPU。

阶段 1：gather_engram_tokens_zos 生成窗口

一个小 kernel，按 (seq) 分配 warp，lane 并行处理 query 维，一次算出 [total_q, N] 的压缩 n-gram 窗口。输入里的 lookup_table 是 GPU 上的压缩映射表——prev 序列的草稿 token 是原始 id，在这里完成压缩。输出与 input_ids 一样是 int32，体积很小（64 × 6 × 4 × 4 B ≈ 6 KB）。

阶段 2：engram_hash 生成索引

另一个小 kernel，对每个 (token, 阶数, 头) 计算哈希并加上表偏移，输出 [2, total_q, 24] 的 int64 全局行号。total_q = 384 时总共 9216 个索引，数据量约 74 KB。images_seq_mask 为空（纯文本请求），图像分支不会触发。

阶段 3：EngramRuntime.fetch ——只是发起

Python 侧调用 self._buffer.fetch(hashes)，得到 {1: hook, 14: hook}。这一步在 CPU 上执行，向通信流提交了 engram_fetch_impl 的发射计划。这里有一个细节需要注意：如果实现把 engram_fetch_impl 放在单独的一条 stream 上，主 stream 需要等待它的完成事件，否则第 1 层的门控可能在数据到齐之前就读了接收区。DeepEP 的用法是把 fetch 的完成信号记在接收区的完成标志上，真正读数的门控 kernel 在发射前等待该标志（或由 stream 依赖保证）。last_gin_requests 与 workspace 里的计数正是服务于这个等待链。

阶段 4：engram_fetch_impl ——网卡开始读

issue warp 按「层 × token tile × 表项」展开，对每个落在本 peer 的索引发出一次 256 字节的 RDMA 读；scale warp 用 LSA 把 8 字节 scale 散写进接收区。__ballot_sync 把请求按 warp 压缩，门铃按阈值聚合。这一阶段与后续 GPU 计算并行。

阶段 5：第 0 层前向

第 0 层没有 Engram，正常执行注意力 + FFN。这几百微秒到毫秒级的时间里，网卡在后台把阶段 4 的读请求落地。第 0 层越慢、读请求越小，隐藏效果越好——这也是为什么 Engram 层放在第 1 层而不是第 0 层：第 0 层至少要提供一段可观的计算窗口。

阶段 6：第 1 层 —— 门控写入

第 1 层进入 hc_post_intervention 时调用 embed_hook()，此时读数据已经完成（或门控 kernel 会等待完成标志）。Engram.forward 做投影、门控、原地写入，然后把结果交给本层的 attn_norm。这一层的 hc_post、Engram、_hc_pre 是三个独立 kernel，不具备 _mega_mhc 融合。

阶段 7：第 2 到 13 层

正常前向。第 14 层再次执行一次门控写入。阶段 4 的第二次 fetch（第 14 层那一份）与阶段 0 到 1 的 fetch 是同时发起的——fetch 一次为所有层发出请求，而不是按层分两次调用。所以第 14 层的延迟窗口从阶段 4 一直持续到阶段 7，比第 1 层宽得多。

阶段 8：采样与 KV 写入

采样得到本步真实 token，写入 KV 池，并把接受数记账给下一拍。压缩历史在 CPU 侧追加本步真实 token。

用一张表回看每一步的量级

CUDA Graph 视角：哪些字段会被「清空重算」

DeepInfer 的 decode 走 CUDA Graph（第十二篇讲了捕获与重放，DeepInfer 源码拆解（十二）：CUDA Graph 捕获与重放）。CudaGraphRunner.capture 用一组填充序列预热，然后逐档捕获模型前向，每一步的输入通过 input_metadata_device.recursive_copy_(input_metadata_source) 原地回填。回填之后有一段「清洁」逻辑：

1
for name in input_metadata_source.get_clean_list():
2
    setattr(input_metadata_source, name, None)
3
    setattr(input_metadata_device, name, None)

get_clean_list 定义了「每次重放前应该被清掉、由 forward 重新生成」的字段：

1
@staticmethod
2
def get_clean_list() -> list[str]:
3
    """Fields to clear (set to None) when reusing a cudagraph; they are recomputed in forward."""
4
    return []
5

6
# InputMetadataV4 覆盖为：
7
return InputMetadata.get_clean_list() + [
8
    "kv_ca_index_len", "kv_swa_index", "kv_ca_index",
9
    "mlp_relu_schedule_metadata", "flash_mla_sched_meta_dict",
10
]

再看 ngram_token_ids 的字段定义：它被声明在 InputMetadata 上，形如 [total_q, max_ngram_size]，在 CP-mirror 下经 gather 变成 [cp_mirror_len, max_ngram_size]。它不在 clean list 里。这意味着两件事：

• 哈希与 fetch 的输入 ngram_token_ids 是图输入的一部分，每步由 prepare_token_ids_zos 写进同一块设备缓冲后再重放，图里记录的指针始终指向这块缓冲；
• 接收区 fetched 同理，必须是一块稳定的、与图生命周期一致的缓冲区，不能在重放之间重新分配。

反过来，如果将来把哈希 kernel 的输出（hashes）也放进 InputMetadata 且被列入 clean list，就会出现「哈希结果在重放前被清空、但图里的 fetch 已经按旧指针发射」的错位。目前源码把 hashes 留在 Python 作用域内、每次 forward 现算现用，正是为了避免这种耦合。读这一块时可以把规则总结成一句话：凡是被 CUDA Graph 记录的输入，必须要么是可原地回填的固定缓冲，要么就在 clean list 里并在 forward 里重建。ngram_token_ids 属于前者，kv_ca_index 之类属于后者。

一个实用的排查方法：如果开启 Engram 后发现每步多出若干次 Graph 重捕获，先看接收区大小与档位是否对齐。接收区按 max_num_decoding_tokens（decode）或 max_num_prefilling_tokens（prefill）分配，而图按 batch_sizes 档位捕获；如果 Env 配置改变了这两个上限，接收区形状会变，图也必须重建。

一个反直觉的顺序

Engram 的投影与门控发生在第 1 层，但投影的权重在模型构造时就准备好了，而输入才在第 1 层才被读取。这意味着如果投影是 FP8 GEMM，GEMM 的 scale 与权重布局必须在捕获前确定——use_tma_aligned_col_major_sf=True 决定了接收区里 scale 的排布，一旦改变这个开关，整张图都要重新捕获。读源码时把「布局开关 → 图捕获」这条因果链记住，能省掉很多配置不一致的调查时间。

边界一：缓存缺失没有「未命中」这个概念

Engram 查表是确定性寻址：哈希命中哪一行就取哪一行，没有「这一行在不在本地」的判断，也没有多级缓存。这和 KV Cache 的前缀缓存、MoE 的专家缓存都不同：

• 前缀缓存命中与否取决于历史请求，命中率随负载变化；
• 专家缓存的命中取决于最近路由的热度；
• Engram 的「命中」在语义上等于「表行存在」，而表是完整的、切好分片的、每 rank 都持有自己那一段。

这意味着首 token 与后续 token 的 Engram 延迟没有「冷热」之分，每个 token 的查表成本几乎恒定（受 RDMA 拥塞与 batch 大小影响）。好处是可预测；代价是不能用「热表」换更快的访问。如果表能放到 GPU 上并利用 HBM 的随机读，单次访问会快得多，但容量不允许。这也是为什么 DEEPINFER_ENGRAM_FORCE_GPU_STORAGE 只适合小模型实验。

边界二：词表压缩改变了「token 是什么」

压缩词表把大小写、重音、空白形式折叠到一起。这对记忆检索是好事，但它也意味着 Engram 看到的 token 序列与主模型 tokenizer 的序列不是一一对应的。调试时必须记住这一点：

• 用原始 tokenizer 打印一句 prompt，数一下 token 数；Engram 的 n-gram 是在压缩空间里数，形状一样但值不同。
• -1 是图像哨兵，不是合法的压缩 id；如果日志里看到 -1 出现在 ngram_token_ids 里，说明该位置是图像或者压缩映射越界（后者通常意味着用原始 id 越过了压缩词表的上界）。
• 压缩映射的构建按原始词表 id 升序遍历，setdefault 分配编号。如果 tokenizer 文件更新（即使内容基本一致，只要 id 顺序有变化），映射就变了，表行也随之全部错位。任何 tokenizer 升级都必须重新训练或重新导出 Engram 表。

边界三：投机解码让 Engram 读到「未来的」草稿

第 1、14 层的门控写入发生在验证之前。被拒绝的草稿对应的 n-gram 会被写入主干，然后被采样后的 KV 回滚「撤销」。这里有一个微妙的问题：主干上的门控写入撤销了，但 KV Cache 里的 K/V 已经被写进去了吗？

时序上，注意力在门控之后执行，被拒绝的草稿对应的 KV 确实会被写入 KV 池（因为模型已经前向过），然后在 update_mtp_runtime / KV 回滚时被截掉。所以对下一拍而言，历史是干净的。但本拍内部，被拒绝草稿的 token 已经在窗口里参与了哈希——这正是第五段分支二所允许的。它带来的误差很小：如果草稿被拒，那几个 n-gram 对应本应不存在的未来，取回的 embedding 被写入主干，然后草稿位被丢弃，不会流入最终的接受序列。代价是每拍多读几行表（nextn × 24 的量级），换来的是查表提前到验证之外。

一个容易踩的坑是：如果实现里把「被拒绝的草稿」也写进 compressed_seq_data，那么下一拍的历史窗口就会包含不存在的 token，哈希在后续所有步都会错位。正确做法是只追加被接受的 token，prev_num_accepted_tokens 正是用来做这件事的唯一依据。

边界四：YOCO、图像与 Engram 的三方约束

• YOCO 截断：Engram 不能落在 consumer 层，否则行布局不一致。
• 图像：图像 token 不参与 n-gram，高阶窗口在图像处截断。
• 二者叠加时，图像 token 只能出现在 full 布局的那一半；如果图像落在 consumer 半区，images_seq_mask 的行索引会与 compact 后的主干错位。源码的断言只拦住了「Engram 在 consumer 层」，没有拦住「图像在 consumer 层」，这是目前版本的一个边界，读代码时值得留意。

边界五：ENP 与 DP×CP 的关系

第十七篇已经讲过 world_size = max(DP×CP, ENP)。对 Engram 来说，这意味着：

• 当 ENP <= DP×CP 时，ENP 组嵌套在 DP 组内，查表流量只在同节点或同 DP 组内绕一圈；
• 当 ENP > DP×CP 时，多出来的 rank 只承载表分片，查表会跨域，RDMA 压力显著上升。

EngramBuffer 的 get_theoretical_config 按 peer 数量估算 QP，如果 ENP 跨域，kNumRDMAPeers 会变大，QP 分配更宽，发射路径也更长。部署时如果发现 Engram 查表成为瓶颈，第一件事是看 ENP 是否被迫跨域。

常见误读一：「Engram 是一层独立的网络」

不是。它是挂在一个 decoder layer 上的干预钩子，共享该层的注意力与前馈。从计算量看，它增加的主要是「查表 + 一次投影 + 一次门控 kernel」，而不是完整的一层。理解成「一层的增强」比理解成「一层」更贴近源码。

常见误读二：「哈希表相当于显存里的一个 embedding」

形状上是 embedding 表（行 × 256 维），但访问方式完全不同：embedding 查表是「一次局部命中」；Engram 查表是「每个 token 24 次远程访问」。两者的成本模型不可互换。把 Engram 表当 embedding 放在 HBM 里，按 V4.1-Flash 的量级根本装不下；把它像 KV 一样放 CPU 上，访问又会变成每步上百 KB 的随机读——所以它需要的是按行哈希确定、可预取、单边读的专用数据面，而不是通用的 embedding 缓存。

常见误读三：「远程取数越早越好」

提前到「请求进入系统」就发起查表听起来更理想，但本实现只在模型前向开始时发起。原因是多步之间 ngram_token_ids 依赖上一步的采样与接受数，再提前就要引入更复杂的跨拍流水。第 1 层到第 14 层之间、第 0 层与第 1 层之间的计算窗口已经是现成的隐藏空间；在 decode 场景下，第 0 层的耗时足以覆盖典型 RDMA 延迟。若把表放到更远的域（例如跨机房），这个窗口可能不够——那时的正确做法不是更早发起，而是把 Engram 层再往后放，或者用分块预取。