Engram in DeepSeek V4.1 Flash: Knowledge Without Recomputing It Every Time

We consider DeepSeek V4.1 Flash a remarkable feat of model engineering. It brings together image and text understanding, sparse experts, compressed sparse attention, and Engram conditional memory, giving knowledge capacity, computation, and runtime resources more distinct roles.

The interesting questions are concrete: which information needs contextual computation, which local patterns can come from learned memory, and which intermediate results can be shared across layers? These decisions determine how much weight data inference reads, how much history it retains, and which work can start early. Implementing the model in an inference engine reveals how those details fit together.

This series follows our actual integration work in zLLM. Our first placement decision was to keep Engram's large tables and projections on the host, while the backbone runs on GPUs. Gating initially ran on the CPU as well; later optimization moved prefill gating onto the GPU.

That decision shaped the rest of the work: weight access, the moment each token enters the memory state, state propagation across GPUs, and whether image positions participate in text lookup all had to become explicit parts of the inference path.

As of September 14, 2026, zLLM has run the complete image-and-text-to-generation path on eight W7900D GPUs using the official safetensors weights. Verified text performance is about 26.06 seconds to the first token for a 50K input, and 20.45 tokens/s for single-request generation. Both figures are from text tests with DSpark disabled: the former is average time to first token for 50,032 input tokens; the latter is average generation speed for 700 output tokens following a 25-token prompt. This first article covers Engram's mechanism and engineering optimizations in full. Other model components and the complete benchmark tables follow in later articles.

Knowledge in memory: Engram

The upper path does not depend on the current layer's hidden state, so lookup and projection can run ahead. The lower path waits for that state before gating. The key/value boxes in both panels refer to the same projection outputs.

Why remember something instead of reconstructing it every time?

Some language-model computation interprets new context, combines information, and reasons. Some reconstructs familiar local patterns, such as representations of common phrases, fixed expressions, and entity names. A conventional Transformer encodes these capabilities together in its weights and reconstructs the relevant representations through network computation for each input.

Engram provides a lookup path for the part suited to memory. During training, reusable local patterns are learned in n-gram memory tables. At inference time, token combinations identify rows, retrieve learned vectors, and let the current context determine how to incorporate them. The official mechanistic analysis suggests this relieves early layers of static-pattern reconstruction, potentially preserving more effective depth for complex reasoning. Official Engram research overview

Consider a familiar proper noun spanning several tokens. A memory table can supply learned features associated with that local combination; subsequent computation determines its role in this particular sentence. This is an illustration of the mechanism, not a claim that each row stores a readable fact or that we have traced a particular phrase to its exact bucket.

Some representations that the network would otherwise reconstruct now have a direct lookup path. More precisely, training assigns some knowledge and patterns to dedicated memory parameters. zLLM loads those weights and executes retrieval; our integration does not extract a separate knowledge base from an older model.

Computational savings and VRAM savings happen at different levels

At the model-design level, Engram assigns some static-pattern reconstruction to conditional memory. The table is large, but each step accesses a fixed number of rows. Adding memory capacity does not require scanning the entire table for every token. This improves the relationship between memory capacity and computation; it does not establish that disabling Engram in the existing V4.1 model would make inference slower, or that enabling it automatically skips Transformer layers.

At the engine level, zLLM uses that sparse access pattern to place the large tables in host memory. VRAM can then hold backbone weights, experts, KV cache, and temporary tensors without also holding the entire Engram table. This saves the capacity that a fully GPU-resident copy of the same table would require. The machine still has to store that memory somewhere.

At approximately 384 million rows × 256 dimensions × 2 layers, the one-byte codes alone occupy about 183 GiB, before scales. This is a calculation from tensor dimensions, not a measured RSS value or measured VRAM reduction. By contrast, retrieving 24 rows in each of two layers requires only 12 KiB of codes per token. Actual operating-system access happens in pages, and projection reads more weights afterward, so 12 KiB is not the total memory traffic of an inference step.

Lookup does not produce a complete answer. Retrieved vectors still pass through WKV projection and contextual gating, and backbone inference continues. Our optimizations therefore address distinct stages: small-row access overhead, projection efficiency, thread and data locality, and early execution of work that does not depend on the hidden state.

Turning knowledge into addressable vectors

Engram adds a memory-read path triggered by the token sequence. Short sequences around the current position determine addresses in a learned table. The current hidden state then determines how much of the retrieved information to inject.

This knowledge exists as vectors in model parameters. There are no directly editable factual entries, and a conversation does not automatically write new facts back into the weights. Input determines where to look; context determines how to use what was retrieved.

Our configuration places Engram at L1 and L14, using zero-based layer numbers. Each position produces 2-grams, 3-grams, and 4-grams, with eight hash heads per length: 24 retrieved rows, each 256 dimensions wide.

For one token at one Engram layer, that is:

3 n-gram lengths × 8 heads × 256 dimensions = 6,144 values

Each layer's table has roughly 384 million rows, but a lookup touches only 24. Large total capacity with sparse per-step access motivated host placement.

Step 1: Get the addresses right

The first requirement is matching the token-to-address mapping. The official implementation normalizes token text and maps tokens with identical normalized forms to a shared compressed ID. Case, accents, and whitespace affect the resulting address; the compressed vocabulary size also affects hash-multiplier generation. Official Engram implementation

zLLM follows the official generation rules, exports the compressed token map offline as engram_token_map.bin, and reads it at runtime. Per-layer multipliers, prime bucket sizes, and offsets are expanded into static tables. The map currently contains 129,280 u32 entries.

The runtime path is straightforward:

token ID
  → compressed token ID
  → current token and the preceding three positions
  → 24 hashes across 2 / 3 / 4-grams
  → 24 row addresses in the corresponding layer's table

An incorrect mapping can still return a correctly shaped vector—it just retrieves different memory. We aligned hashing before validating projection and gating.

Step 2: Read selected rows and compute the projection on the host

The weight layer's engram_embedding_rows reads only selected rows. The official tables use MXFP8: E4M3 codes with E8M0 scales in groups of 32 values. The corresponding scales are read with the codes, and the selected rows are decoded on the CPU.

The initial integration read each row at a file offset. Hot data could be served by the operating-system page cache, while cold accesses could reach storage. Later optimization switched to read-only mmap and page-by-page warmup during loading, described below. Warmup is not page locking: memory pressure can still evict pages.

Projection weights are explicitly held in memory. The retrieved 6,144-dimensional vector passes through wkv to produce the key and value used for gating. Each layer's matrix has shape 25,600 × 6,144; converting it to BF16 during preparation takes roughly 300 MiB. The initial implementation used AVX2/FMA with Rayon row parallelism. Later work replaced that with AVX-512 BF16 and persistent worker teams.

Reading only 24 table rows therefore does not make Engram free. Projection still reads a substantial matrix. Keeping it on the CPU saves VRAM but puts pressure on host memory bandwidth. Cold reads, scheduling, and CPU/GPU synchronization also affect latency and must be measured separately.

Step 3: Inject memory at the correct layer boundary

Engram modifies the hidden state at the entrance to L1 and L14, while it still contains the multiple streams of the expanded mHC representation.

The CPU implementation computes key and value, then gates each hidden-state stream separately. Hidden state and key have their own normalization factors. After computing their correlation score dot, it applies:

gate = sigmoid(copysign(sqrt(max(abs(dot), 1e-6)), dot))
hidden += gate × value

The gate controls how much memory the current context absorbs. Our implementation stores the elementwise product of the q and k weights in advance; normalization and gate computation remain separate for each stream.

The initial layer-entry hook invoked CPU Engram and returned the updated hidden state to the GPU backbone. The optimized implementation separates lookup, projection, and gating: tables and projection stay on the CPU, precomputed prefill results go to GPU gating, and decode can retain CPU gating. Placement depends on the execution path.

Step 4: Keep memory history local to the session

Each token must enter the hash history exactly once. Both Engram layers read the history at that same position using their own hash parameters. Appending the token again at each layer would shift subsequent n-grams.

zLLM therefore separates appending tokens from applying Engram. Prefill appends a sequence; decode appends one token at a time. The L1 and L14 hooks use the positions already established.

Weights can be shared, but sequence histories must remain isolated. The implementation shares projection weights through Arc, creates fresh hash state when forking a session, and clears history on reset. Session-isolation fixes were necessary for serving the model correctly.

Images introduce another boundary. Image spans insert DEAD markers to prevent text n-grams from crossing those positions, and image rows skip Engram gating. We completed both behaviors during multimodal numerical validation so image placeholders would not be treated as ordinary text-memory inputs.

Read optimization: from per-row pread to mmap warmup

The first row reader was sufficient to establish correctness. But each row contains only 256 quantized values, and separately reading codes and scales creates many small file operations. Prefill repeats this across the whole input, making fixed per-call costs increasingly visible.

Optimization commit e9b539e4 added read-only mmap to the safetensors path. Inference copies selected rows from the mapped region, reducing per-row pread system calls. The large table remains quantized; it is not expanded wholesale into F32.

Loading also warms the Engram codes and scales: it advises sequential access and readahead, touches each page, then restores MADV_RANDOM for runtime lookup. This moves some first-touch faults and storage reads into loading. The costs are startup time and host-memory occupancy; the benefit needs to be measured after warmup.

This path does not use mlock. Mapping a file alone does not ensure its pages are in RAM; explicit page touches perform the warmup. Memory pressure and page residency still matter when describing the running system.

AVX-512 BF16: reuse loaded weights across tokens

The initial bottleneck: repeated WKV scans

Initially, each input token required a separate GEMV. BF16 already reduced weight bytes, but a token-by-token prefill loop still repeatedly scanned the same approximately 300 MiB matrix.

The main optimization was to put multiple tokens into one matrix multiplication and reuse loaded weights across those inputs. AVX-512 BF16 provides the instructions; batching and layout determine how effectively they work.

Pack BF16 pairs across 16 outputs

The optimized WKV layout is:

[output block][input-dimension pair][16 output channels]

Each u32 packs two BF16 weights. The inner loop loads weight pairs for 16 output channels, broadcasts a token's two input values, and calls _mm512_dpbf16_ps to accumulate into 16 F32 outputs. Multiple tokens keep independent accumulators while sharing that weight load.

The inspected version handles input tiles and tails in groups of 30, 16, 8, 4, 2, and 1 rows. A 30-row tile can reuse one load across as many as 30 tokens. Single-row decode does not receive that batching benefit.

Match input layout to the inner loop

Inputs use a pair-major layout, [input-dimension pair][token row]. Values for the same dimension pair across tokens sit next to each other, avoiding large token-stride jumps. Official MXFP8 table rows can also decode directly into this packed BF16 input, avoiding a full intermediate F32 expansion followed by conversion back to BF16.

This change exposed a real layout bug. Each token's input contains 24 heads, and packing must preserve the head-to-pair mapping. Changing the packer without checking the consuming indices can produce valid dimensions with incorrectly ordered content. We added comparisons between packed MXFP8 and the F32 decoding path, and between batched BF16 computation and a scalar reference.

Runtime checks for avx512f and avx512bf16 select the supported kernel. Changes to input precision handling and accumulation require numerical validation, not just a faster microbenchmark. Local reference tests, real-input generation, and end-to-end performance are separate checks.

NUMA placement: put threads and weights together

Projection reads large weight regions continuously, so the relationship between a thread's CPU and the memory node holding its pages matters. If a worker runs on one socket while much of its data resides on another, remote access can offset vectorization gains.

Identify the physical cores available to the process

The implementation first calls sched_getaffinity and considers only allowed CPUs. It then reads physical_package_id and core_id from sysfs, groups by physical package, and removes duplicate SMT siblings. With two packages, L1 and L14 use different CPU groups.

The code groups by package. That matches the placement strategy for this dual-socket machine, but it is not a complete NUMA-topology solution for arbitrary hardware: one socket can expose several NUMA nodes. Moving to another machine requires checking its topology rather than assuming socket and NUMA node always mean the same thing.

Persistent workers with fixed partitions

Each layer creates a persistent EngramTeam. Workers bind to selected CPUs at startup and process fixed output-block ranges. They park when idle, wake for work, and notify the submitter when finished. This stabilizes thread placement and work partitioning, reducing the migration and locality changes possible with general-purpose dynamic scheduling.

The selection leaves cores available for other work. When a package offers at least 16 physical cores, Engram uses three quarters of them. A fully available 32-core package therefore supplies 24 workers, leaving scheduling headroom for GPU submission, transfers, and serving. This is not a cpuset guarantee that those other threads run exclusively on the remaining cores.

Make loading and first-touch follow computation

The fixed workers repack WKV directly into its final buffer. Initial writes therefore occur on the CPUs that will later consume those partitions. Table warmup also temporarily binds to a CPU from the corresponding layer's group before touching mapped pages.

This combines thread affinity with first-touch placement. The implementation does not use mbind to forcibly migrate existing pages. If another process has already loaded file pages, or the system applies a different memory policy, touching pages from a bound thread does not guarantee relocation. Actual page distribution and remote accesses still need inspection.

AVX-512 improves computation after data arrives. NUMA placement determines where those bytes come from. Leaving CPU headroom helps projection avoid starving GPU submission. All three serve the same complete pipeline.

Upload near the GPU that consumes the projection

After moving prefill gating onto the GPU, projection results still needed an upload. Measurements found about 15 ms for L1 and 90 ms for an equally sized L14 result: the L14 CPU team and its consuming GPU were on different NUMA nodes.

We added pinned staging near the destination GPU so the transfer could use a local page-locked buffer. In that round of adjacent experiments, 50K time to first token fell from roughly 26.67–26.84 seconds to 25.87–25.88 seconds, with matching output hashes. Those figures describe that experiment; later measurements provide the final benchmark table.

Binding also needs measurement. Moving all GPU submission threads to the other NUMA node produced 28.47 seconds and was reverted. Successful data placement and upload staging do not imply that binding every thread according to PCIe topology will always help.

Run projection early: overlap CPU work with the GPU

Engram addresses depend only on token history, and WKV projection depends only on retrieved vectors. Only gating needs the current layer's hidden state. That dependency allows the substantial computation to start early.

The optimized path prepares both layers' Engram batches at stage 0 and launches projection tasks. GPUs advance through the backbone concurrently, then obtain the results at L1 or L14. If projection is complete, only the remaining work is exposed at layer entry; otherwise, execution still waits. L14 has more preceding backbone work and therefore a larger potential overlap window. Profiling must establish how much latency is actually hidden.

For precomputed prefill, projected key/value results go to the GPU for gating, keeping the large hidden-state batch on the device. Subsequent decode commit 153598f6 also starts projection early but retains CPU gating. This reduces exposed waits and some transfers while leaving the tables and WKV computation on the CPU.

In the adjacent decode A/B test, a fixed 25-token prompt and 700-token output improved from an average 19.417 to 20.454 tokens/s, a 5.34% gain, with the complete output-text hash unchanged. The full benchmark table follows in the performance article. AVX-512, NUMA placement, warmup, and pipeline changes all contributed to prefill optimization; the overall gain cannot be attributed to one change, and batching gains do not directly translate into single-token decode gains.

Coming next

We plan to continue with one article per day on the remaining engineering work:

• Part 2: From weight loading to an eight-GPU text pipeline
• Part 3: Adding vision—finding differences through intermediate tensors
• Part 4: Benchmarks—prefill, decode, and before-and-after comparisons