Prerequisite reading: This document walks through every model file insrc/hymo/models/line by line. Start here if you want to understand exactly how a forward pass flows through the 32-layer hybrid stack.
Files covered:
-model.py— the top-levelHyMomodel andbuild_hymofactory
-gdn.py— Gated Delta Net linear-attention block
-gdn_triton.py— Custom Triton kernel for the GDN recurrence
-mla.py— Multi-Head Latent Attention with MQA-4 grouping
-moe.py— DeepSeekMoE, DenseFFN, SwiGLU, and aux-loss-free routing
-mtp.py— Multi-Token Prediction heads (depth=2)
-rope.py— Rotary Position Embedding (partial RoPE)
Stale-section notice (as of commitaf89c48): Sections 5.5.3 ("Step 5: Gated Delta Rule Recurrence", lines ~1655–1664) and 6 (the Triton kernel section, ~line 1790) describe a hypotheticalfla-first / Triton-second fallback that was never shipped. The shipped code path is Triton-only (gdn_triton.py::triton_gated_delta_rule) with nofladependency, and the eager path (gdn.py::_gated_delta_rule) is the reference for tests. The current source-of-truth for the kernel isoptimization.md§3 andkernels.md. Treat thefla/chunk_gated_delta_rulesnippets in those sections as historical plan-time documentation that illustrates the recurrence interface, not as the implementation.
Companion concepts (theory tier — read these for the why behind each section; section numbers refer to this walkthrough's table of contents):
-gdn-and-mla.md— §1 high-level architecture
-model-architecture.md— §4 MLA
-gdn-and-mla.md— §5–§6 GDN + Triton kernel
-model-architecture.md— §7 RoPE
-gdn-and-mla.md— §8 MoE
-gdn-and-mla.md— §9 MTP
-optimization.md— §10 μP
-kernels.md— §6 Triton kernel
See../README.mdfor the full reading order and cross-reference table.
Table of Contents#
- High-Level Architecture
- Configuration System (core/config.py)
- The HyMo Stack (model.py)
- Multi-Head Latent Attention (mla.py)
- Gated Delta Net Linear Attention (gdn.py)
- Custom Triton GDN Kernel (gdn_triton.py)
- Rotary Position Embedding (rope.py)
- Mixture of Experts (moe.py)
- Multi-Token Prediction (mtp.py)
- Initialization — status note
- End-to-End Forward Pass Trace
- Parameter Count Breakdown
1. High-Level Architecture — Design Rationale & Stack Design#
*This section explains why HyMo is built the way it is.* The 3:1 hybrid ratio, the MLA placement strategy, the asymmetric FFN design, and the parameter sharing scheme all have specific mathematical and empirical motivations.
1.1 The Hybrid Attention Hypothesis#
1.1.1 Why Not Pure Full Attention?#
Standard transformer attention (MHA/GQA) has two costs that scale with sequence length T:
| Cost | Complexity | Bottleneck |
|---|---|---|
| Compute (QK^T) | O(T² × d) | FLOPs grow quadratically |
| Memory (KV cache) | O(T × n_h × h_d) | Cache grows linearly |
At T=4096, d=896, n_heads=16, head_dim=128:
- Compute: 4096² × 896 ≈ 15 billion FLOPs per layer
- Memory: 4096 × 16 × 128 × 2 bytes ≈ 16 MB per layer
For 8 MLA layers: 128 MB of KV cache. This is manageable at T=4096 but becomes prohibitive at T=32K (1 GB) or T=128K (4 GB).
1.1.2 Why Not Pure Linear Attention?#
Linear attention (GDN) reduces both costs:
| Cost | Complexity | Savings vs MHA |
|---|---|---|
| Compute | O(T × S × D) | S=32, D=32 → ~400× faster |
| Memory | O(S × D) per layer | Fixed state, regardless of T |
But linear attention has a fundamental limitation: limited long-range discrimination. The state H_t ∈ R^{S×D} is a compressed summary — it can store information but cannot selectively attend to arbitrary past positions.
1.1.3 The 3:1 Ratio: Cheap Recurrence + Precise Retrieval#
The key insight: most language processing is local and sequential (grammar, syntax, nearby context), while occasional long-range retrieval is needed (coreference, distant dependencies).
┌──────────────────────────────────────────────────────────────┐
│ Token Processing Pipeline: │
│ │
│ GDN layers (24×): O(T) recurrence — cheap, local context │
│ ↓ │
│ MLA layers (8×): O(T²) attention — expensive, precise │
│ ↓ │
│ GDN layers (24×): O(T) recurrence — cheap, local context │
│ ↓ │
│ MLA layers (8×): O(T²) attention — expensive, precise │
│ ↓ │
│ ... (repeats every 4 layers) │
└──────────────────────────────────────────────────────────────┘Why 3:1 specifically?
The ratio is derived from empirical observations in hybrid attention literature (Jamba, Zamba, Zamba2, StripedHyena):
| Ratio | Quality (PPL) | Throughput | Notes |
|---|---|---|---|
| 1:1 (16 MLA + 16 GDN) | Best | Slowest | Full attention everywhere |
| 3:1 (8 MLA + 24 GDN) | ~0.02 PPL loss | 3-5× faster | HyMo's choice |
| 7:1 (4 MLA + 28 GDN) | ~0.1 PPL loss | 5-7× faster | Too few attention layers |
| All GDN | Significant loss | Fastest | Cannot do precise retrieval |
The 3:1 ratio captures ~95% of full-attention quality at ~40% of the compute cost.
1.2 MLA Placement Strategy#
1.2.1 Why Every 4th Layer?#
MLA blocks are placed at positions {0, 4, 8, 12, 16, 20, 24, 28} — every 4th layer.
Design rationale:
- Uniform coverage: With 32 layers, 8 MLA blocks spaced every 4 layers ensures no token goes more than 3 layers without a full-attention retrieval step
- GDN state refresh: After 3 GDN layers, the recurrent state may drift. MLA provides a "checkpoint" that re-anchors the representation
- Gradient flow: MLA's quadratic attention provides a high-bandwidth gradient pathway back through the stack, preventing gradient vanishing in the deep recurrence
Layer: 0 1 2 3 4 5 6 7 8 ...
Type: MLA GDN GDN GDN MLA GDN GDN GDN MLA ...
↓ ↓ ↓ ↓ ↓
Retrieve Accumulate Retrieve Accumulate Retrieve
precise local precise local precise
context features context features context1.2.2 Why MLA at Position 0?#
Meta FAIR (Bae et al., 2025) recommends "Never place Transformer blocks at the front" — front placement leads to performance degradation. However, HyMo places MLA at position 0 for two reasons:
- MLA's latent bottleneck acts as a learned summarization — the
kv_lora_rank=128bottleneck compresses input embeddings into a structured representation, which is beneficial even at the input stage - Empirical validation: The first 1k training steps will validate this choice. If loss spikes at position 0, MLA moves to position 1
Open question for v1.1: If MLA-at-0 hurts, swap to GDN-at-0 (position 1 becomes MLA).
1.3 Asymmetric Feed-Forward: MoE on MLA, Dense on GDN#
1.3.1 Why Not MoE Everywhere?#
The intuition: MoE dispatch overhead is amortized on expensive layers.
| Layer Type | Per-Token Cost | MoE Overhead | Overhead Fraction |
|---|---|---|---|
| MLA | ~5.8M params (high) | ~9M params | ~15% |
| GDN | ~2.5M params (low) | ~9M params | ~36% |
Putting MoE on GDN layers wastes 36% of the GDN compute budget on routing. On MLA, the overhead is only 15%.
1.3.2 Routing Noise Compounding#
More critically: GDN has recurrent state. If a token is routed to the wrong expert in a GDN block:
GDN (with MoE):
t=0: token routed to expert 3 (wrong) → corrupted value → stored in H_0
t=1: H_1 = exp(g·A)·H_0 + ... → corrupted H_0 propagates to H_1
t=2: H_2 = exp(g·A)·H_1 + ... → corruption compounds
...The routing error persists in the recurrent state and corrupts all future predictions.
MLA (with MoE):
MLA (with MoE):
t=0: token routed to expert 3 (wrong) → corrupted attention output
t=1: MLA recomputes from scratch (no recurrence) → error is localMLA's attention is stateless — routing errors are local to one position and corrected at the next layer.
1.3.3 No FFN on GDN#
GDN blocks are recurrence-only — the gated delta rule output goes straight through out_proj to the residual stream. There is no SwiGLU FFN after the recurrence:
- No extra dense parameters on the 24 GDN layers
- Every GDN parameter contributes to the recurrence (in_proj, conv1d, b/c/dt/g_proj, out_proj, skip_proj)
- Compute is spent only on the recurrence itself
1.4 Stack Configuration Summary#
┌─────────────────────────────────────────────────────────────┐
│ HyMo v1.0 Architecture │
├─────────────────────────────────────────────────────────────┤
│ Total layers: 32 │
│ MLA layers: 8 (positions 0, 4, 8, ..., 28) │
│ GDN layers: 24 (all other positions) │
│ MLA:GDN ratio: 1:3 (8:24) │
│ │
│ Model dim (d): 896 │
│ Heads (MLA): 16 │
│ Head dim (MLA): 128 (32 RoPE + 96 NoPE) │
│ KV groups (MLA): 4 (MQA-4) │
│ GDN heads: 40 (d_inner=1280, headdim=32) │
│ GDN state: S=32, D=32 │
│ │
│ MoE (MLA only): 16 routed + 1 shared, top-2 │
│ FFN (GDN only): Dense SwiGLU, inter_dim=2560 │
│ │
│ Vocab: 64,256 (BPE-64k + 256-byte fallback) │
│ Embedding: Weight-tied with output head │
│ Logit softcap: 15.0 (tanh stabilization) │
│ MTP depth: 2 (predict tokens t+2, t+3) │
│ MTP weights: [0.3, 0.1] │
│ │
│ Active params: ~434M │
│ Stored params: ~1.13B (MoE experts are stored) │
│ Training tokens: 30B (~69× params-in-tokens) │
│ Context length: 4,096 │
└─────────────────────────────────────────────────────────────┘1.5 The Five Invariants#
| # | Invariant | Why It Matters |
|---|---|---|
| 1 | Asymmetric FFN | MoE on MLA only — avoids routing noise in recurrence |
| 2 | Partial-RoPE 25% | Position/content decoupling — cleaner KV compression |
| 3 | MQA-4 | 4 KV groups — compression/quality sweet spot |
| 4 | FP32 master weights | Numerical stability in BF16 training |
| 5 | Weight-tied embeddings | 57.6M params saved — same tensor for input/output |
1.6 Visual Stack Diagram#
Input Tokens (B, T)
│
▼
┌─────────────────────────────────────────────────────────────┐
│ Embedding: (B, T) → (B, T, 896) │
│ Weight: 64,256 × 896 = 57.6M params (tied with output head) │
└─────────────────────────────────────────────────────────────┘
│
▼
┌─────── MLA block 0 (layer_idx=0) ──────────────────────────┐
│ Pre-norm → Multi-Head Latent Attention (MQA-4) │
│ • Query: 896 → 224 (bottleneck) → 2048 (16 heads × 128) │
│ • KV: 896 → 128 (latent) + 32 (RoPE key) │
│ • Decoupled RoPE on 32/128 dims (25%) │
│ Pre-norm → DeepSeekMoE (16 routed + 1 shared, top-2) │
│ • SwiGLU experts: 896 → 2304 → 896 │
│ • Aux-loss-free load balancing (EMA bias adjustment) │
└─────────────────────────────────────────────────────────────┘
│
▼
┌─────── GDN block 1 (layer_idx=1) ──────────────────────────┐
│ In-proj: 896 → 1280 │
│ Depthwise Conv1d (kernel=4) → SiLU │
│ Gated Delta Rule: H_t = exp(g·A)·H_{t-1} + b⊗v │
│ • 40 heads, S=32, D=32, O(T) recurrence │
│ • RoPE on full value (32/32 dims = 100%) │
│ Skip connection + gate │
│ Out-proj: 1280 → 896 + residual │
│ Dense SwiGLU FFN: 896 → 2560 → 896 │
└─────────────────────────────────────────────────────────────┘
│
▼
┌─────── GDN block 2 (layer_idx=2) ──────────────────────────┐
│ (Same as block 1) │
└─────────────────────────────────────────────────────────────┘
│
▼
┌─────── GDN block 3 (layer_idx=3) ──────────────────────────┐
│ (Same as block 1) │
└─────────────────────────────────────────────────────────────┘
│
▼
┌─────── MLA block 4 (layer_idx=4) ──────────────────────────┐
│ (Same as block 0) │
└─────────────────────────────────────────────────────────────┘
│
▼
... (repeats: MLA at 8, 12, 16, 20, 24, 28; GDN elsewhere)
│
▼
┌─────────────────────────────────────────────────────────────┐
│ Final RMSNorm → (B, T, 896) │
└─────────────────────────────────────────────────────────────┘
│
▼
┌─────────────────────────────────────────────────────────────┐
│ Output Head: Linear(896 → 64,256) — weight-tied with embed │
│ Logit Softcap: 15.0 × tanh(logits / 15.0) │
└─────────────────────────────────────────────────────────────┘
│
▼
Logits (B, T, 64,256) + MTP auxiliary logits (depth=2)2. Configuration System (core/config.py) — Design Philosophy & Mathematical Derivations#
Why a config system matters: Every architectural knob in HyMo is defined in ModelConfig (a frozen dataclass). This section explains the mathematical relationships between config fields, how configs are derived, and why certain validation rules exist.
2.1 The No-Hardcoding Principle#
Every hyperparameter lives in configs/hymo_750m.yaml. The only numeric literals allowed in code are architectural constants (e.g., vocab_size = 64_256 for the BPE-64k + 256-byte tokenizer).
Why? Because hyperparameters are tuned iteratively. If head_dim=128 is hardcoded in model code, changing it requires editing multiple files. With config, a single YAML change propagates everywhere.
2.2 Config Flow: YAML → Model#
from hymo import load_config, build_hymo
config = load_config("configs/hymo_750m.yaml") # → HyMoConfig
model = build_hymo(config) # → HyMo(config.model)┌─────────────────────────────────────────────────────────────┐
│ YAML File (hymo_750m.yaml) │
│ ├── model: ModelConfig │
│ ├── optimizer: OptimizerConfig │
│ ├── scheduler: SchedulerConfig │
│ ├── training: TrainingConfig │
│ └── run: RunConfig │
└─────────────────────────────────────────────────────────────┘
│
▼ load_config()
┌─────────────────────────────────────────────────────────────┐
│ HyMoConfig (top-level aggregator) │
│ ├── .model → ModelConfig │
│ ├── .optimizer → OptimizerConfig │
│ ├── .scheduler → SchedulerConfig │
│ ├── .training → TrainingConfig │
│ └── .run → RunConfig │
└─────────────────────────────────────────────────────────────┘
│
▼ build_hymo(config)
┌─────────────────────────────────────────────────────────────┐
│ HyMo(config.model) │
│ ├── embed = Embedding(64256, 896) │
│ ├── head = Linear(896, 64256) ← tied to embed │
│ ├── layers[0] = MLABlock(config, 0) │
│ ├── layers[1] = GatedDeltaNetBlock(config, 1) │
│ ├── ... │
│ └── _mtp = MultiTokenPrediction(config, self) │
└─────────────────────────────────────────────────────────────┘2.3 The Five Sub-Configs#
| Config class | YAML group | Key Fields | Purpose |
|---|---|---|---|
ModelConfig | model | dim, n_heads, n_layers, mtp_depth | Architecture knobs |
OptimizerConfig | optimizer | lr, betas, weight_decay, muon_lr | Optimizer hyperparams |
SchedulerConfig | scheduler | warmup_frac, stable_frac, decay_frac | WSD schedule |
TrainingConfig | training | batch_size, seq_len, grad_accum_steps | Training dynamics |
RunConfig | run | name, seed, output_dir | Experiment metadata |
2.4 Frozen Dataclasses: Why Immutable?#
@dataclass(frozen=True)
class ModelConfig:
dim: int = 896
n_heads: int = 16
# ... 30+ fieldsWhy frozen? The training loop is a complex state machine. If config is mutable, a bug could accidentally modify dim mid-training:
# BUG: accidentally mutate config during training
config.dim = 1024 # ← FrozenInstanceError raised immediatelyHow to derive variants:
# Correct: create a new config with modified field
from dataclasses import replace
small_config = replace(config, dim=512, n_heads=8)The derive_config() helper computes dependent fields automatically:
# derive_config() computes:
# - qk_rope_head_dim from head_dim and rope_ratio
# - qk_nope_head_dim from head_dim and rope_ratio
# - kv_lora_rank from empirical formula
# - q_lora_rank from empirical formula2.5 Mathematical Relationships Between Config Fields#
2.5.1 Dimension Arithmetic#
d = 896 # model dimension
n_heads = 16 # number of attention heads
head_dim = 128 # per-head dimension
qk_rope_head_dim = 32 # RoPE subspace (25% of head_dim)
qk_nope_head_dim = 96 # Content subspace (75% of head_dim)
Invariants:
head_dim = qk_rope_head_dim + qk_nope_head_dim # 128 = 32 + 96
q_lora_rank = 224 # Query bottleneck
kv_lora_rank = 128 # KV bottleneck
v_head_dim = 128 # Value dimension2.5.2 GDN Dimension Arithmetic#
gdn_d_inner = 1280 # GDN working dimension
gdn_headdim = 32 # GDN per-head dimension
gdn_d_state = 32 # State matrix dimension
gdn_d_conv = 4 # Causal convolution kernel size
gdn_n_heads = gdn_d_inner / gdn_headdim = 1280 / 32 = 40 # GDN heads2.5.3 MoE Dimension Arithmetic#
n_routed = 16 # Number of routed experts
n_shared = 1 # Number of shared experts
n_activated = 2 # Top-k routing
moe_inter_dim = 2304 # Expert FFN intermediate dimension
capacity_factor = 1.5 # Max tokens per expert
capacity = capacity_factor × (batch_tokens × n_activated) / n_routed
= 1.5 × (4 × 4096 × 2) / 16
= 1.5 × 512
= 768 tokens per expert2.6 Validation Rules#
Every config validates on construction. Key rules:
class ModelConfig:
def __post_init__(self):
# Rule 1: Head dimension consistency
assert self.qk_rope_head_dim + self.qk_nope_head_dim == self.head_dim
# 32 + 96 = 128 ✓
# Rule 2: GQA compatibility
assert self.n_heads % self.n_kv_groups == 0
# 16 % 4 = 0 ✓ (4 query heads per KV group)
# Rule 3: MTP weight count
assert len(self.mtp_loss_weights) == self.mtp_depth
# [0.3, 0.1] has 2 elements = mtp_depth=2 ✓
# Rule 4: Layer count
assert self.n_layers % 4 == 0
# 32 % 4 = 0 ✓ (MLA every 4 layers)
# Rule 5: MLA position count
assert len(self.mla_positions) == self.n_layers // 4
# 8 MLA positions = 32 / 4 ✓Cross-field validation in config_validation.py:
- VRAM budget check: total params × bytes + optimizer states + activations ≤ A100 80GB
- Throughput check: tokens_per_step / step_time ≥ target throughput
- Data check:
total_tokens / batch_size≥ number of unique documents
2.7 Computed Properties#
Config fields that are derived, not specified:
class ModelConfig:
@property
def mla_positions(self) -> tuple[int, ...]:
"""MLA at positions 0, 4, 8, ..., n_layers-4."""
return tuple(range(0, self.n_layers, 4))
@property
def gdn_positions(self) -> tuple[int, ...]:
"""GDN at all non-MLA positions."""
return tuple(i for i in range(self.n_layers) if i not in self.mla_positions)
@property
def nope_hybrid_gdn_positions(self) -> tuple[int, ...]:
"""GDN positions immediately after MLA (when NoPE enabled)."""
mla_set = set(self.mla_positions)
return tuple(
i for i in range(self.n_layers)
if i not in mla_set and (i - 1) in mla_set
)
# Returns {3, 7, 11, 15, 19, 23, 27} when enabled2.8 The hymo_750m.yaml Structure#
model:
dim: 896
n_heads: 16
n_layers: 32
vocab_size: 64256
max_seq_len: 4096
# MLA-specific
q_lora_rank: 224
kv_lora_rank: 128
qk_rope_head_dim: 32
qk_nope_head_dim: 96
v_head_dim: 128
n_kv_groups: 4
# GDN-specific
gdn_d_inner: 1280
gdn_d_state: 32
gdn_d_conv: 4
gdn_headdim: 32
# MoE-specific
n_routed: 16
n_shared: 1
n_activated: 2
moe_inter_dim: 2304
capacity_factor: 1.5
# MTP-specific
mtp_depth: 2
mtp_loss_weights: [0.3, 0.1]
logit_softcap: 15.0
optimizer:
lr: 3.0e-4
muon_lr: 0.02
betas: [0.9, 0.95]
weight_decay: 0.1
scheduler:
warmup_frac: 0.02
stable_frac: 0.85
decay_frac: 0.13
min_lr_ratio: 0.05
training:
batch_size: 4
seq_len: 4096
grad_accum_steps: 32
bf16: true
grad_clip: 1.03. The HyMo Stack (model.py) — Assembly, Embeddings & Forward Flow#
This section explains how the 32-layer stack is assembled, how embeddings work, and the complete forward pass flow. The design decisions around weight tying, gradient checkpointing, and logit stabilization are covered in depth.
3.1 Why Weight-Tied Embeddings?#
3.1.1 The Mathematical Argument#
Standard transformer:
Embedding: V × d = 64,256 × 896 = 57.5M parameters
Output head: d × V = 896 × 64,256 = 57.5M parameters
Total: 115M parameters (embedding alone = 6.2% of model)Weight-tied transformer:
Embedding: V × d = 57.5M parameters
Output head: same tensor (0 additional parameters)
Total: 57.5M parameters (saved 57.5M)Why does this work? The output logits are:
logits = h · W^T where h ∈ R^d (hidden state), W ∈ R^{V×d} (embedding)This is the dot product between the hidden state and every embedding vector. The intuition: a token's embedding is its "identity" — predicting the next token means finding which embedding is most similar to the current hidden state.
3.1.2 When Weight Tying Breaks Down#
Weight tying assumes the input embedding space and the output logit space are isomorphic. This works well for:
- Language models (next-token prediction)
- Models where the vocabulary is shared between input and output
It may not work for:
- Multilingual models with different input/output vocabularies
- Models with separate encoder/decoder vocabularies
HyMo uses weight tying because it's a decoder-only LM with a shared BPE vocabulary.
3.2 Embedding: The First Transformation#
self.embed = nn.Embedding(config.vocab_size, config.dim)
# Weight: (64,256, 896) — each token has a 896-dim vectorForward:
x = self.embed(tokens) # (B, T) → (B, T, 896)The embedding is a lookup table: for each integer token ID, return the corresponding row from the weight matrix. This is equivalent to a one-hot encoding followed by a linear projection:
embed(t) = one_hot(t) @ W_embedBut the lookup is O(1) per token, not O(V × d).
μP initialization: Embedding weights use std = 1/√d = 1/√896 ≈ 0.033 (not 1/d like other parameters). This ensures the embedding vectors have unit variance at initialization, which is important for the first token to have reasonable magnitude.
3.3 Layer Assembly: The Loop#
self.layers = nn.ModuleList()
for i in range(config.n_layers): # 32 iterations
if i in mla_positions: # {0, 4, 8, ..., 28}
self.layers.append(MLABlock(config, layer_idx=i))
else:
use_rope = i not in nope_hybrid # True for all 24 GDN in v1.0
self.layers.append(
GatedDeltaNetBlock(config, layer_idx=i, use_rope=use_rope)
)Key design decisions:
nn.ModuleList— not a plain list. This ensures all layer parameters are registered for:- Optimizer parameter groups
- FSDP sharding
model.parameters()iteration
layer_idxpassed to each block — used for:- Positional encoding offsets (if needed)
- Per-layer configuration (e.g.,
use_ropeflag) - Debugging and logging
use_ropeflag — In v1.0, all 24 GDN layers use RoPE (use_rope=True). In v1.1, 7 GDN layers at positions {3, 7, 11, 15, 19, 23, 27} may getuse_rope=False(NoPE-hybrid ablation).
3.4 Gradient Checkpointing: Memory vs Compute Trade-off#
def _run_layers(self, x: torch.Tensor) -> torch.Tensor:
for layer in self.layers:
use_cp = getattr(layer, "use_checkpoint", False)
if use_cp and self.training:
x = checkpoint(layer, x, use_reentrant=False)
else:
x = layer(x)
return xWhy checkpointing?
During training, activations (intermediate tensors) must be stored for backward pass. For 32 layers, this is:
Activation memory = 32 layers × (B × T × d) × bytes_per_element
= 32 × 4 × 4096 × 896 × 2 bytes
= ~940 MB (BF16)Gradient checkpointing reduces this by:
- Only storing activations at "checkpoints" (layer boundaries)
- Recomputing activations during backward pass
- Saves ~60% activation memory at ~33% compute overhead
# Without checkpointing:
# Forward: store all 32 layer activations → 940 MB
# Backward: use stored activations → 0 MB recomputation
# With checkpointing every 4 layers:
# Forward: store 8 checkpoint activations → 235 MB
# Backward: recompute 3 layers between checkpoints → 235 MB recomputation
# Total: 235 MB (saved 705 MB) at ~33% more computeIn v1.0, neither MLABlock nor GatedDeltaNetBlock enable checkpointing by default. It's reserved for the production run if activation memory becomes tight.
3.5 Logit Softcapping: Preventing Pathological Logits#
def softcap(self, logits):
if self.logit_softcap <= 0:
return logits
return self.logit_softcap * torch.tanh(logits / self.logit_softcap)The problem: During long training runs, logits can grow unboundedly. This causes:
- Numerical instability in softmax (overflow)
- Gradient explosion
- Training divergence
The solution: Tanh softcapping maps logits to [-softcap, softcap]:
softcapped = softcap × tanh(logits / softcap)Mathematical properties:
softcapped ≈ logitswhen|logits| ≪ softcap(linear regime)softcapped → ±softcapwhen|logits| ≫ softcap(saturation)- Gradient:
d(softcapped)/d(logits) = 1 - tanh²(logits/softcap)(vanishes at saturation)
At softcap=15.0:
- Logits in
[-5, 5]: ~95% of the range, gradient ~0.91 - Logits in
[-10, 10]: ~99.8% of the range, gradient ~0.64 - Logits beyond ±15: gradient → 0, capped at ±15
This is a PaLM-style stabilization technique that prevents training divergence without losing representational capacity.
3.6 Factory Function: build_hymo#
def build_hymo(config: HyMoConfig) -> HyMo:
return HyMo(config.model)Why a factory? The separation between HyMo(config) and build_hymo(config) exists so that future versions can perform post-initialization steps:
def build_hymo(config: HyMoConfig) -> HyMo:
model = HyMo(config.model)
# No custom init pass — PyTorch module defaults + inline gate/GDN init
# (see optimization.md). FSDP wrapping happens in the trainer.
return modelCurrently build_hymo is a pass-through, but the indirection is a deliberate architectural choice for extensibility.
3.7 Complete Forward Flow#
# HyMo.forward(tokens)
def forward(self, tokens: torch.Tensor) -> torch.Tensor:
hidden = self.forward_with_hidden(tokens)[1]
logits = self.head(hidden)
return self.softcap(logits)
# HyMo.forward_with_hidden(tokens, start_pos=0)
def forward_with_hidden(self, tokens, start_pos=0):
x = self.embed(tokens) # (B, T) → (B, T, 896)
x = self._run_layers(x) # 32-layer stack
hidden = self.norm(x) # RMSNorm → (B, T, 896)
logits = self.head(hidden) # (B, T, 64,256)
return self.softcap(logits), hiddenData flow:
tokens: (B, T) = (4, 4096) int64
│
▼ embed()
x: (B, T, 896) = (4, 4096, 896) BF16
│
▼ _run_layers() — 32 iterations
│ ├── MLABlock(0): x = x + attn(attn_norm(x))
│ │ x = x + moe(moe_norm(x))
│ ├── GatedDeltaNetBlock(1): x = GDN(x) + FFN(x)
│ ├── GatedDeltaNetBlock(2): x = GDN(x) + FFN(x)
│ ├── GatedDeltaNetBlock(3): x = GDN(x) + FFN(x)
│ ├── MLABlock(4): (same as block 0)
│ ├── ... (28 more layers)
│ └── GatedDeltaNetBlock(31): last GDN layer
│
x: (4, 4096, 896) BF16
│
▼ norm() — RMSNorm
hidden: (4, 4096, 896) BF16
│
▼ head() — Linear(896 → 64,256), weight = embed.weight.T
logits: (4, 4096, 64,256) BF16
│
▼ softcap() — 15.0 × tanh(logits / 15.0)
output: (4, 4096, 64,256) BF163.8 Code Walkthrough (src/hymo/models/model.py)#
"""HyMo: 32-layer hybrid attention model (architecture doc §2)."""
from __future__ import annotations
from typing import cast
import torch
from torch import nn
from torch.nn import functional as F
from torch.utils.checkpoint import checkpoint
from hymo.core.config import HyMoConfig, ModelConfig
from hymo.models.gdn import GatedDeltaNetBlock
from hymo.models.mla import MLABlock
from hymo.models.moe import DeepSeekMoE
from hymo.models.mtp import MultiTokenPrediction
from hymo.models.registry import MODELS
@MODELS.register("hymo")
class HyMo(nn.Module):
"""32-layer hybrid: 24 GDN + 8 MLA with 3:1 ratio."""
def __init__(self, config: ModelConfig) -> None:
super().__init__()
self._config = config
# Embedding (weight-tied with output head)
self.embed = nn.Embedding(config.vocab_size, config.dim)
if config.tie_embeddings:
self.head = nn.Linear(config.dim, config.vocab_size, bias=False)
self.head.weight = self.embed.weight # SAME tensor
else:
self.head = nn.Linear(config.dim, config.vocab_size, bias=False)
self.norm = nn.RMSNorm(config.dim)
self.logit_softcap = config.logit_softcap
# Layer assembly
mla_positions = config.mla_positions
nope_hybrid = config.nope_hybrid_gdn_positions
self.layers = nn.ModuleList()
for i in range(config.n_layers):
if i in mla_positions:
self.layers.append(MLABlock(config, layer_idx=i))
else:
use_rope = i not in nope_hybrid
self.layers.append(
GatedDeltaNetBlock(config, layer_idx=i, use_rope=use_rope)
)
# MTP heads
if config.mtp_depth > 0:
self._mtp = MultiTokenPrediction(config, main_model=self)
else:
self._mtp = None
def forward(self, tokens: torch.Tensor) -> torch.Tensor:
hidden = self.forward_with_hidden(tokens)[1]
logits = self.head(hidden)
return self.softcap(logits)
def forward_with_hidden(self, tokens: torch.Tensor, start_pos: int = 0):
x = self.embed(tokens)
x = self._run_layers(x)
hidden = self.norm(x)
logits = self.head(hidden)
return self.softcap(logits), hidden
def _run_layers(self, x: torch.Tensor) -> torch.Tensor:
for layer in self.layers:
use_cp = getattr(layer, "use_checkpoint", False)
if use_cp and self.training:
x = checkpoint(layer, x, use_reentrant=False)
else:
x = layer(x)
return x
def softcap(self, logits: torch.Tensor) -> torch.Tensor:
if self.logit_softcap <= 0:
return logits
return self.logit_softcap * torch.tanh(logits / self.logit_softcap)
def build_hymo(config: HyMoConfig) -> HyMo:
"""Build HyMo from config. Factory for future extensibility."""
return HyMo(config.model)4. Multi-Head Latent Attention (mla.py) — Mathematical Foundations & Design Decisions#
File: src/hymo/models/mla.py
This section builds MLA from first principles. If you want to understand why each architectural choice exists — the mathematical motivation, the information-theoretic justification, the empirical evidence from DeepSeek-V2/V3, and the trade-offs — read this section carefully. The code walkthrough follows after the conceptual foundation.
4.1 The Problem: Why Standard Attention Doesn't Scale#
4.1.1 KV Cache Memory Bottleneck#
In standard Multi-Head Attention (MHA), for a sequence of length T, model dimension d, n_heads heads, and head_dim = d / n_heads:
KV cache per layer = 2 × n_heads × head_dim × T × sizeof(dtype)At HyMo's configuration (n_heads=16, head_dim=128, T=4096, BF16):
= 2 × 16 × 128 × 4096 × 2 bytes = 33,554,432 bytes ≈ 32 MB per layerWith 8 MLA layers: 256 MB total KV cache — just for the attention state. This doesn't include activations, gradients, or optimizer states.
For inference at longer contexts (e.g., T=32K), this becomes 2 GB — a major bottleneck for serving.
4.1.2 Quadratic Complexity#
Standard attention computes:
Attention(Q, K, V) = softmax(QK^T / √d) VThe QK^T matrix is (T × T), giving O(T²) time and memory complexity. At T=4096, that's ~16M pairwise interactions per head per layer. For 16 heads across 8 layers: ~2 billion operations per forward pass.
4.1.3 The Low-Rank Hypothesis (DeepSeek-V2 Key Finding)#
DeepSeek-V2 (Dai et al., 2024) observed empirically that the KV activation matrix is approximately low-rank. Specifically, for a layer's key/value projections:
K ∈ R^{T × (n_heads × head_dim)}, V ∈ R^{T × (n_heads × head_dim)}The effective rank of K and V is much smaller than their full dimension. This means most of the information in K and V can be compressed into a much smaller latent space with minimal loss.
Intuition: Language has strong local and hierarchical structure. Tokens that are close in context share similar representations. The "degrees of freedom" in the KV space are far fewer than the raw dimension suggests.
4.2 MLA: Low-Rank KV Compression from First Principles#
4.2.1 Mathematical Formulation#
MLA factorizes the KV projection into a low-rank bottleneck:
Standard MHA: K = X W_K, V = X W_V where W_K, W_V ∈ R^{d × (n_heads × head_dim)}
MLA: K = X W_kv_a W_kv_b, V = X W_kv_a W_kv_bBut with a critical modification: joint compression with decoupled RoPE.
Let d = 896 (model dim), kv_lora_rank = 128 (latent dim), n_kv_groups = 4 (MQA groups).
wkv_a: R^d → R^{kv_lora_rank + qk_rope_head_dim} # 896 → 128 + 32 = 160
└─> kv_latent (128) + k_pe (32) # Split after projection
kv_norm: RMSNorm on kv_latent # Stabilize latent scale
wkv_b: R^{kv_lora_rank} → R^{n_kv_groups × (qk_nope + v_head_dim)} # 128 → 4 × (96 + 128) = 896
└─> k_nope (96 per group) + v (128 per group)Why this factorization?
| Aspect | Standard MHA | MLA |
|---|---|---|
| KV projection params | 2 × d × (n_h × h_d) = 1.84M | d × 160 + 128 × 896 = 257K |
| KV cache per token | 2 × n_h × h_d = 4,096 | kv_lora_rank + n_kv_groups × qk_rope = 128 + 128 = 256 |
| Compression ratio | 1× | 16× |
The 16× cache reduction comes from storing only the 128-dim latent + 32-dim RoPE key per token, instead of full 2,048-dim K and V.
4.2.2 Decoupled RoPE: Why Separate Position from Content?#
Standard RoPE applies rotation to the full key/query:
RoPE(q) = q ⊙ cos + rotate(q) ⊙ sinProblem: The rotation mixes position and content information. For low-rank compression, this is suboptimal — the position information "pollutes" the content space that we're trying to compress.
MLA's solution: Decouple the position encoding:
- Content path:
kv_latent(128-dim) — pure content, no position info, compressed - Position path:
k_pe(32-dim) — pure position, shared across all KV groups, uncompressed
K_full = concat(k_nope, k_pe_rotated) # 96 + 32 = 128 per groupThis means:
- Content compression is clean —
kv_latentcontains only semantic information - Position is explicit —
k_peis a dedicated position encoding subspace - Shared computation — One
k_peper token serves all 4 KV groups (broadcast)
Design Decision (from DeepSeek-V2): The decoupled RoPE is applied to a fixed 25% of head_dim (32 of 128). This 25% ratio was empirically found optimal — large enough to encode position, small enough to leave 75% for content.
4.2.3 MQA-4: Multi-Query Attention with 4 Groups#
Standard attention variants:
| Variant | KV Heads | Query Heads | Ratio | Notes |
|---|---|---|---|---|
| MHA | 16 | 16 | 1:1 | Full capacity, max memory |
| GQA | 8 | 16 | 1:2 | Good balance |
| MQA-4 | 4 | 16 | 1:4 | MLA choice |
| MQA | 1 | 16 | 1:16 | Max compression, potential quality loss |
Why MQA-4 for MLA?
- Compression synergy: MLA already compresses KV to latent. Fewer KV groups = less decompression work
- Hardware efficiency: 4 groups maps well to GPU warp/wavefront boundaries
- Quality preservation: 4 groups is the "sweet spot" from DeepSeek-V3 ablation — 1 group (MQA) loses too much capacity, 8 groups (GQA) defeats the compression purpose
Each KV group serves n_heads / n_kv_groups = 16/4 = 4 query heads. This is implemented via Grouped Query Attention (GQA) in SDPA:
- Q: (B, 16, T, 128)
- K: (B, 4, T, 128) → broadcast to 16 heads internally
- V: (B, 4, T, 128) → broadcast to 16 heads internally
4.3 Query Path: Low-Rank with Per-Head Scaling#
4.3.1 Two-Stage Projection#
wq_a: R^d → R^{q_lora_rank} # 896 → 224 (bottleneck)
q_norm: RMSNorm on q_lora_rank # Per-token normalization
wq_b: R^{q_lora_rank} → R^{n_heads × head_dim} # 224 → 2,048Why two stages?
- Parameter efficiency:
896×224 + 224×2048 = 659Kvs direct896×2048 = 1.84M(2.8× savings) - The bottleneck forces the model to learn a compressed query representation
RMSNormafter bottleneck stabilizes the scale before expansion
4.3.2 Per-Dimension Learnable Scaling (q_norm_qk, k_norm_qk)#
After projecting to per-head space, MLA applies element-wise learnable scales:
q = q.view(B, T, H, head_dim) * q_norm_qk.view(1, 1, H, head_dim)
# q_norm_qk ∈ R^{n_heads × head_dim} = R^{2048}And similarly for K's RoPE portion:
k_pe_normed = k_pe * k_norm_qk.view(1, 1, G, D_pe)
# k_norm_qk ∈ R^{n_kv_groups × qk_rope_head_dim} = R^{128}Why per-dimension scaling?
- Attention logit control: The softmax in attention is sensitive to logit magnitude. These scales let the model learn which dimensions should contribute more/less to attention scores
- Stabilization: Initialized to 1 (identity), they provide a "temperature" per dimension that can adapt during training
- DeepSeek finding: This simple trick consistently improves training stability and final perplexity across scales
Analogy: Think of this as a learned "attention head importance weighting" but at the dimension level rather than head level.
4.4 Complete Forward Pass: Mathematical Trace#
Let's trace through with concrete shapes at HyMo config:
B=4, T=4096, d=896H=16, G=4, D_pe=32, D_nope=96, D_v=128q_lora=224, kv_lora=128
Step 1: Query Computation#
x: (B, T, d) = (4, 4096, 896)
q_latent = RMSNorm(wq_a(x)) # (4, 4096, 224)
q = wq_b(q_latent) # (4, 4096, 2048)
q = q.view(4, 4096, 16, 128) # (B, T, H, head_dim)
q = q * q_norm_qk # Element-wise scale (4, 4096, 16, 128)
q_pe = q[..., :32] # (4, 4096, 16, 32) — position dims
q_nope = q[..., 32:] # (4, 4096, 16, 96) — content dims
q_pe_rot = RoPE(q_pe) # (4, 4096, 16, 32) — apply rotationStep 2: KV Compression#
kv = wkv_a(x) # (4, 4096, 160)
kv_latent = kv[..., :128] # (4, 4096, 128) — compressed content
k_pe = kv[..., 128:] # (4, 4096, 32) — position key
kv_latent = RMSNorm(kv_latent) # (4, 4096, 128)
kv_out = wkv_b(kv_latent) # (4, 4096, 896)
kv_out = kv_out.view(4, 4096, 4, 224) # (B, T, G, D_nope + D_v)
k_nope = kv_out[..., :96] # (4, 4096, 4, 96)
v = kv_out[..., 96:] # (4, 4096, 4, 128)Step 3: Position Key Processing (Decoupled RoPE)#
# Broadcast k_pe to all 4 KV groups with per-group scaling
k_pe_normed = k_pe.unsqueeze(2) * k_norm_qk # (4, 4096, 1, 32) * (1, 1, 4, 32)
# → (4, 4096, 4, 32) via broadcast
k_pe_rot = RoPE(k_pe_normed) # (4, 4096, 4, 32)Step 4: Assemble for SDPA#
q_assembled = concat(q_pe_rot, q_nope, dim=-1) # (4, 4096, 16, 128)
k_assembled = concat(k_pe_rot, k_nope, dim=-1) # (4, 4096, 4, 128)
# SDPA expects (B, n_heads, T, head_dim)
q_sdpa = q_assembled.permute(0, 2, 1, 3) # (4, 16, 4096, 128)
k_sdpa = k_assembled.permute(0, 2, 1, 3) # (4, 4, 4096, 128)
v_sdpa = v.permute(0, 2, 1, 3) # (4, 4, 4096, 128)Step 5: Grouped Query Attention#
# Native GQA (PyTorch 2.5+) handles 4 KV heads → 16 Q heads internally
out = scaled_dot_product_attention(q_sdpa, k_sdpa, v_sdpa, enable_gqa=True)
# out: (4, 16, 4096, 128)Step 6: Output Projection#
out = out.permute(0, 2, 1, 3).contiguous().view(4, 4096, 2048)
y = wo(out) # (4, 4096, 896)4.5 Design Decision Deep Dives#
4.5.1 Why kv_lora_rank = 128?#
The latent dimension is the information bottleneck. Too small → information loss. Too large → defeats compression.
DeepSeek-V2/V3 ablation results:
| kv_lora_rank | KV Cache (MB/layer) | PPL (validation) |
|---|---|---|
| 64 | 16 | +0.15 |
| 128 | 32 | baseline |
| 256 | 64 | -0.02 |
| 512 | 128 | -0.01 |
128 is the knee of the curve — diminishing returns beyond this point. HyMo uses 128.
4.5.2 Why q_lora_rank = 224?#
Query compression is less aggressive because:
- Queries are not cached (only K,V are)
- Queries need to discriminate across all positions — more capacity helps
- 224 = 22.5% of full query dim (2048) — enough for rich query representations
4.5.3 Why Partial RoPE (25%)?#
Full RoPE on all dimensions would:
- Make KV compression harder (position + content mixed)
- Waste capacity on position encoding
No RoPE would:
- Break position awareness for long-range dependencies
25% (32/128) is the empirically optimal trade-off from DeepSeek-V2:
- Sufficient for position encoding (32 dims can encode 4096+ positions via RoPE)
- Leaves 75% for content discrimination
4.5.4 Why RMSNorm on Latents?#
q_latent = RMSNorm(wq_a(x))
kv_latent = RMSNorm(kv_latent)RMSNorm (Root Mean Square Normalization) on the bottleneck representations:
- Prevents scale explosion through the bottleneck
- Stabilizes gradients — the bottleneck is a gradient highway
- No learnable parameters (unlike LayerNorm) — just per-feature scaling
4.6 MLA Block: Attention + MoE with Pre-Norm Residuals#
The MLABlock wraps the attention with a DeepSeekMoE feed-forward:
class MLABlock(nn.Module):
def __init__(self, config, layer_idx=0):
self.attn_norm = RMSNorm(config.dim) # Pre-norm for attention
self.attn = MultiHeadLatentAttention(config, layer_idx)
self.moe_norm = RMSNorm(config.dim) # Pre-norm for MoE
self.moe = DeepSeekMoE(config, layer_idx) # 16 experts + 1 shared
def forward(self, x):
x = x + self.attn(self.attn_norm(x)) # Attention residual
x = x + self.moe(self.moe_norm(x)) # MoE residual
return xWhy MoE only on MLA blocks?
| Reason | Explanation |
|---|---|
| Routing overhead | MoE dispatch is expensive. MLA is already compute-heavy — amortizes the overhead |
| Routing noise | GDN has recurrent state; routing errors compound. MLA is stateless — errors are local |
| Capacity allocation | MLA layers need more FFN capacity for precise retrieval; GDN does cheap aggregation |
4.7 Code Walkthrough (src/hymo/models/mla.py)#
Now with the mathematical foundation, let's walk the actual implementation:
class MultiHeadLatentAttention(nn.Module):
"""Multi-Head Latent Attention (MLA) with MQA-4 grouping.
Features decoupled RoPE on 25% of the head dimension and low-rank compression:
- Query: x -> wq_a -> RMSNorm -> wq_b -> q (split into q_nope and q_pe).
- KV: x -> wkv_a -> (kv_latent, k_pe), and kv_latent -> RMSNorm -> wkv_b -> (k_nope, v).
"""
def __init__(self, config: ModelConfig, layer_idx: int = 0) -> None:
super().__init__()
self.layer_idx = layer_idx
self._config = config
d = config.dim
n_heads = config.n_heads
n_kv_groups = config.n_kv_groups
qk_rope_head_dim = config.qk_rope_head_dim
qk_nope_head_dim = config.qk_nope_head_dim
kv_lora_rank = config.kv_lora_rank
q_lora_rank = config.q_lora_rank
v_head_dim = config.v_head_dim
# Query path projection: two-stage low-rank
self.wq_a = nn.Linear(d, q_lora_rank, bias=False)
self.q_norm = nn.RMSNorm(q_lora_rank)
self.wq_b = nn.Linear(
q_lora_rank, n_heads * (qk_rope_head_dim + qk_nope_head_dim), bias=False
)
# KV compression: joint latent + decoupled RoPE
self.wkv_a = nn.Linear(d, kv_lora_rank + qk_rope_head_dim, bias=False)
self.kv_norm = nn.RMSNorm(kv_lora_rank)
self.wkv_b = nn.Linear(
kv_lora_rank,
n_kv_groups * (qk_nope_head_dim + v_head_dim),
bias=False,
)
# Output projection
self.wo = nn.Linear(n_heads * v_head_dim, d, bias=False)
# Per-head RMSNorm scaling applied post-split
self.q_norm_qk = nn.Parameter(
torch.ones(n_heads * (qk_rope_head_dim + qk_nope_head_dim))
)
self.k_norm_qk = nn.Parameter(
torch.ones(n_kv_groups * qk_rope_head_dim)
)
# Decoupled Rotary Position Embedding
self.rope = RotaryEmbedding(
head_dim=qk_rope_head_dim,
max_seq_len=config.max_seq_len,
theta=config.rope_theta,
)
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""Forward pass for MLA mapping input (B, T, d) -> (B, T, d)."""
B, T, _ = x.shape
H, G = self.n_heads, self.n_kv_groups
D_pe, D_nope, D_v = (
self.qk_rope_head_dim, self.qk_nope_head_dim, self.v_head_dim
)
# Query path: compress -> scale -> split to decoupled RoPE / nope
q_latent = self.q_norm(self.wq_a(x))
q = self.wq_b(q_latent)
head_dim_q = D_pe + D_nope
q = q.view(B, T, H, head_dim_q) * self.q_norm_qk.view(1, 1, H, head_dim_q)
q_pe = q[..., :D_pe]
q_nope = q[..., D_pe:]
q_pe_rot = self.rope.apply_rope(
q_pe.permute(0, 2, 1, 3)
).permute(0, 2, 1, 3)
# KV path: compress to joint latent & shared key position (k_pe)
kv = self.wkv_a(x)
kv_latent, k_pe = kv.split(
[self._config.kv_lora_rank, D_pe], dim=-1
)
kv_latent = self.kv_norm(kv_latent)
kv_out = self.wkv_b(kv_latent).view(B, T, G, D_nope + D_v)
k_nope = kv_out[..., :D_nope]
v = kv_out[..., D_nope:]
# Scale and apply RoPE to the shared position key k_pe across groups
k_pe_normed = k_pe.unsqueeze(2) * self.k_norm_qk.view(1, 1, G, D_pe)
k_pe_rot = self.rope.apply_rope(
k_pe_normed.permute(0, 2, 1, 3)
).permute(0, 2, 1, 3)
# Reconstruct keys and queries by concatenating position-encoded and non-position-encoded components
q_assembled = torch.cat([q_pe_rot, q_nope], dim=-1)
k_assembled = torch.cat([k_pe_rot, k_nope], dim=-1)
# SDPA with MQA-4 broadcast. Prefers hardware GQA support (PyTorch >= 2.5), otherwise falls back.
q_sdpa = q_assembled.permute(0, 2, 1, 3)
k_sdpa = k_assembled.permute(0, 2, 1, 3)
v_sdpa = v.permute(0, 2, 1, 3)
heads_per_group = H // G
try:
out = F.scaled_dot_product_attention(
q_sdpa, k_sdpa, v_sdpa,
enable_gqa=True,
)
except (TypeError, RuntimeError):
k_sdpa = k_sdpa.repeat_interleave(heads_per_group, dim=1)
v_sdpa = v_sdpa.repeat_interleave(heads_per_group, dim=1)
out = F.scaled_dot_product_attention(q_sdpa, k_sdpa, v_sdpa)
out = out.permute(0, 2, 1, 3).contiguous().view(B, T, H * D_v)
y = self.wo(out)
return cast(torch.Tensor, y)
class MLABlock(nn.Module):
"""MLA Block combining latent attention, MoE feed-forward, pre-norms, and residuals."""
def __init__(self, config: ModelConfig, layer_idx: int = 0) -> None:
super().__init__()
self.layer_idx = layer_idx
self._config = config
from hymo.models.moe import DeepSeekMoE
self.attn_norm = nn.RMSNorm(config.dim)
self.attn = MultiHeadLatentAttention(config, layer_idx=layer_idx)
self.moe_norm = nn.RMSNorm(config.dim)
self.moe = DeepSeekMoE(config, layer_idx=layer_idx)
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""Forward pass applying pre-norm Attention followed by pre-norm MoE."""
x = x + self.attn(self.attn_norm(x))
x = x + self.moe(self.moe_norm(x))
return x4.8 Summary: MLA Configuration in HyMo v1.0#
| Parameter | Value | Rationale |
|---|---|---|
n_heads | 16 | Standard for 896-dim |
n_kv_groups | 4 | MQA-4: compression/quality sweet spot |
head_dim | 128 | 896/16 = 56? No — expanded to 128 for capacity |
qk_rope_head_dim | 32 | 25% partial RoPE |
qk_nope_head_dim | 96 | 75% content |
v_head_dim | 128 | Full value dimension |
q_lora_rank | 224 | Query bottleneck (2.8× param savings) |
kv_lora_rank | 128 | KV bottleneck (16× cache reduction) |
rope_theta | 10,000 | Standard RoPE base |
max_seq_len | 4,096 | Training context |
5. Gated Delta Net Linear Attention (gdn.py) — Mathematical Foundations#
File: src/hymo/models/gdn.py
This section builds GDN from first principles. The gated delta rule is the core innovation that makes linear attention competitive with full attention. We derive the recurrence, explain the information-theoretic properties, and show why the read/write key decoupling is strictly more expressive than standard linear attention.
5.1 The Problem: Linear Attention vs Quadratic Attention#
5.1.1 Standard (Quadratic) Attention#
Standard self-attention computes:
Attention(Q, K, V) = softmax(QK^T / √d) VComplexity: O(T² × d) compute, O(T × d) memory for QK^T
Expressiveness: Can attend to any past position with arbitrary weights.
Limitation: Quadratic scaling makes long contexts expensive.
5.1.2 Linear Attention#
Linear attention replaces the softmax with a kernel trick:
LinearAttn(Q, K, V) = ϕ(Q) (ϕ(K)^T V)Where ϕ is a kernel feature map. This can be computed as:
S_t = S_{t-1} + ϕ(k_t) v_t^T ← State update
o_t = ϕ(q_t)^T S_t ← OutputComplexity: O(T × S²) where S is the state dimension (S=32 in HyMo)
Expressiveness: Limited by the state dimension S. Cannot selectively attend to arbitrary past positions.
5.1.3 The Key Limitation of Linear Attention#
Linear attention's state S_t ∈ R^{S×S} is a fixed-size summary of all past tokens. Once information is compressed into S, it cannot be selectively retrieved — the read is a fixed linear combination.
Analogy: Linear attention is like writing notes in a notebook (state), then summarizing each page. Standard attention is like having the full book available to search.
5.2 The Gated Delta Rule: A Better Recurrence#
5.2.1 The Recurrence#
For each head at position t:
H_t = exp(g_t · A) ⊙ H_{t-1} + b_t ⊗ v_t ← State update (delta rule)
o_t = c_t^T · H_t ← Read from stateWhere:
H_t ∈ R^{S × D}— hidden state matrix (S=d_state=32, D=headdim=32)b_t ∈ R^S— write key (what to write to state)c_t ∈ R^S— read key (what to read from state)v_t ∈ R^D— value vectorg_t ∈ R— per-head scalar gate (output of sigmoid(dt_proj))A ∈ R^S— negative decay vector (-exp(A_log), making it strictly negative)exp(g_t · A) ⊙— element-wise decay: each state dimension decays independently
5.2.2 Why "Delta Rule"?#
The name comes from the delta rule in online learning:
Δw = η × (target - prediction) × inputIn GDN:
ΔH = b_t ⊗ v_t ← outer product of write key and valueThis is the "delta" — the change to the state matrix. The write key b_t determines where in the state to write, and the value v_t determines what to write.
5.2.3 Why Read/Write Decoupling Matters#
Standard linear attention:
S_t = S_{t-1} + k_t v_t^T ← same k for read and write
o_t = q_t^T S_t ← same q for read and writeGated delta net:
H_t = α H_{t-1} + b_t ⊗ v_t ← b for write
o_t = c_t^T H_t ← c for readThe key difference: b and c are independent learned projections. This means:
- Write with one pattern, read with another: The model can learn to write information using one representation and retrieve it using a completely different representation
- Selective erasure: The gate
g_tcan selectively decay specific state dimensions while preserving others - Expressiveness: This is strictly more expressive than standard linear attention (where b=c=k)
Formal argument: Standard linear attention is a special case of GDN whereb = c = kandg = 1(no gating). GDN generalizes this by allowing independent read/write keys and learned gating.
5.3 Mathematical Properties#
5.3.1 Decay and Forgetting#
The decay term exp(g_t · A) provides exponential forgetting:
H_t = exp(g_t · A) ⊙ H_{t-1} + b_t ⊗ v_tUnrolling the recurrence:
H_t = exp(g_t · A) ⊙ exp(g_{t-1} · A) ⊙ ... ⊙ exp(g_1 · A) ⊙ H_0
+ exp(g_t · A) ⊙ exp(g_{t-1} · A) ⊙ ... ⊙ exp(g_2 · A) ⊙ b_1 ⊗ v_1
+ ...
+ exp(g_t · A) ⊙ b_{t-1} ⊗ v_{t-1}
+ b_t ⊗ v_tThe contribution of token at position k < t to the current state is:
weight(k, t) = ∏_{j=k+1}^{t} exp(g_j · A)
= exp(Σ_{j=k+1}^{t} g_j · A)This is a learned exponential decay — the model controls how fast past information fades.
5.3.2 Multi-Scale Memory#
The A_log parameter is initialized with multi-scale decay:
a_init = log(arange(1, n_heads + 1)) # [log(1), log(2), ..., log(40)]
# Head 0: A = -1, exp(-1) = 0.37 (remembers ~3 steps)
# Head 10: A = -10, exp(-10) ≈ 0.000045 (forgets almost instantly)
# Head 39: A = -39, exp(-39) ≈ 10^{-17} (remembers only current step)Why multi-scale? Different heads specialize for different timescales:
- Slow heads (small |A|): Remember distant context (coreference, topic)
- Fast heads (large |A|): Focus on immediate context (syntax, local semantics)
5.3.3 State Capacity#
The state H_t ∈ R^{S×D} has S×D = 32×32 = 1,024 parameters per head.
Information-theoretic bound: The state can store at most 1,024 independent scalar values. With 40 heads, total state capacity is 40 × 1,024 = 40,960 scalars.
Compare to full attention: At T=4096, full attention stores 4,096 × 128 = 524,288 scalars per head (QK^T matrix). GDN's 1,024 is ~500× smaller.
The trade-off: GDN is much faster but stores less information. The 3:1 ratio ensures MLA provides the "full book" access when needed.
5.4 The GDN Block Architecture#
5.4.1 Input Projection#
self.in_proj = nn.Linear(d_model, d_inner, bias=False) # 896 → 1280
self.conv1d = nn.Conv1d(d_inner, d_inner, d_conv, groups=d_inner,
padding=d_conv - 1, bias=False) # Depthwise causal convWhy 1280? The GDN paper recommends d_inner / d_model ≈ 1.43. At d_model=896:
d_inner = 896 × 1.43 ≈ 1280This ensures the GDN path is not under-provisioned relative to MLA and MoE.
Why depthwise conv1d? Provides local context before the global recurrence. Each channel gets its own 1D convolution (no cross-channel mixing), adding ~0 parameters but capturing 4-token local patterns.
5.4.2 The Projections#
self.b_proj = nn.Linear(d_inner, n_heads * d_state, bias=False) # 1280 → 40×32 = 1280
self.c_proj = nn.Linear(d_inner, n_heads * d_state, bias=False) # 1280 → 1280
self.dt_proj = nn.Linear(d_inner, n_heads, bias=False) # 1280 → 40
self.g_proj = nn.Linear(d_inner, d_inner, bias=False) # 1280 → 1280
self.out_proj = nn.Linear(d_inner, d_model, bias=False) # 1280 → 896
self.skip_proj = nn.Linear(d_model, d_model, bias=False) # 896 → 896Why separate b_proj and c_proj? The decoupled read/write keys are the core innovation. If they were the same (like standard linear attention), we'd lose the expressiveness advantage.
Why g_proj? The gate g_t controls the decay rate. It's computed from the input via a learned projection, then passed through sigmoid to get a value in [0, 1].
5.4.3 The Gate Computation#
g_in = self.g_proj(v) # (B,T,1280)
g = F.sigmoid(g_in) # (B,T,1280) — element-wise in [0,1]
g_gate = g.view(B, T, H, D).mean(dim=-1) # (B,T,40) — scalar per headWhy mean over D? The gate is per-head, not per-dimension. Averaging over the D=32 dimensions collapses the per-head gate to a single scalar.
Why sigmoid? The gate modulates the decay rate:
g → 0: Slow decay (remember more)g → 1: Fast decay (forget more)g = 0.5: Moderate decay
The sigmoid ensures the gate is in [0, 1], which is important for training stability.
5.5 Complete Forward Pass Trace#
Let's trace with concrete shapes at HyMo config:
B=4, T=4096, d_model=896, d_inner=1280, S=32, D=32, H=40
Step 1: Input Projection#
x: (4, 4096, 896)
v_in = in_proj(x) # (4, 4096, 1280)Step 2: Causal Depthwise Conv1d#
v_conv = conv1d(v_in.transpose(1, 2)) # (4, 1280, 4096+3)
v_conv = v_conv[:, :, :4096] # (4, 1280, 4096) — trim padding
v = SiLU(v_conv.transpose(1, 2)) # (4, 4096, 1280)Step 3: Compute Projections#
b_in = b_proj(v) # (4, 4096, 1280) → reshape to (4, 4096, 40, 32)
c_in = c_proj(v) # (4, 4096, 1280) → reshape to (4, 4096, 40, 32)
dt_in = dt_proj(v) # (4, 4096, 1280) → (4, 4096, 40)
g_in = g_proj(v) # (4, 4096, 1280)
g = sigmoid(g_in) # (4, 4096, 1280)
g_gate = g.view(4, 4096, 40, 32).mean(dim=-1) # (4, 4096, 40)Step 4: RoPE on Value#
v = v.view(4, 4096, 40, 32) # (4, 4096, 40, 32)
# RoPE applied to full 32 dims (rope_dim == D)
v_rot = RoPE(v) # (4, 4096, 40, 32)Step 5: Gated Delta Rule Recurrence#
> **HISTORICAL — never shipped; the repo uses the hand-written Triton kernel (gdn_triton.py).**
> The `fla`-first fallback below is plan-time fiction; the shipped path is
> `triton_gated_delta_rule` (gdn_triton.py), with `_gated_delta_rule` (gdn.py)
> as the eager reference for tests.
# Priority 1: FLA kernel
try:
from fla.layers.gated_delta_net import chunk_gated_delta_rule
o = chunk_gated_delta_rule(c, b, v, A_log, dt_in, g_gate, chunk_size)
# Priority 2: Triton kernel
except ImportError:
from hymo.models.gdn_triton import triton_gated_delta_rule
o = triton_gated_delta_rule(v, b, c, g_gate, A_log)
# Priority 3: PyTorch fallback
except ImportError:
o = self._gated_delta_rule(v, b, c, g_gate)The recurrence (PyTorch fallback):
h = zeros(B, H, S, D, dtype=float32) # State: (4, 40, 32, 32)
for t in range(T):
alpha = exp(g_t * A) # Decay factor: (4, 40, 32)
h = alpha.unsqueeze(-1) * h + b_t.unsqueeze(-1) * v_t.unsqueeze(-2)
o_t = einsum("bhs,bhsd->bhd", c_t, h)Step 6: Skip Connection and Gating#
o = o + D.view(1, 1, H, 1) * v # Skip: (4, 4096, 40, 32)
o = o * g_gate.unsqueeze(-1) # Gate: (4, 4096, 40, 32)Step 7: Output Projection#
o_flat = o.view(4, 4096, 1280) # (4, 4096, 1280)
y = out_proj(o_flat) + skip_proj(x) # (4, 4096, 896) + residual5.6 The Skip Connection and D Parameter#
self.D = nn.Parameter(torch.ones(n_heads)) # (40,) — initialized to 1The skip connection adds the original value directly to the output:
o_t = c_t^T H_t + D · v_tWhy? This ensures the block can pass information directly without going through the state. The D parameter controls how much of the direct signal to include.
Initialization to 1: At initialization, the block is approximately an identity function:
o_t ≈ v_t (since H_0 = 0 and D = 1)This is a residual-like initialization that ensures stable training at the start.
5.7 Code Walkthrough (src/hymo/models/gdn.py)#
class GatedDeltaNetBlock(nn.Module):
"""Gated Delta Net block with O(T) recurrence."""
def __init__(self, config: ModelConfig, layer_idx: int, use_rope: bool = True) -> None:
super().__init__()
self._config = config
self.layer_idx = layer_idx
d_model = config.dim # 896
d_inner = config.gdn_d_inner # 1280
d_state = config.gdn_d_state # 32
d_conv = config.gdn_d_conv # 4
headdim = config.gdn_headdim # 32
n_heads = d_inner // headdim # 40
self.d_inner = d_inner
self.n_heads = n_heads
self.headdim = headdim
self.d_state = d_state
# Projections
self.in_proj = nn.Linear(d_model, d_inner, bias=False)
self.conv1d = nn.Conv1d(d_inner, d_inner, d_conv, groups=d_inner,
padding=d_conv - 1, bias=False)
self.b_proj = nn.Linear(d_inner, n_heads * d_state, bias=False)
self.c_proj = nn.Linear(d_inner, n_heads * d_state, bias=False)
self.dt_proj = nn.Linear(d_inner, n_heads, bias=False)
self.g_proj = nn.Linear(d_inner, d_inner, bias=False)
self.out_proj = nn.Linear(d_inner, d_model, bias=False)
self.skip_proj = nn.Linear(d_model, d_model, bias=False)
# Learnable decay and skip
a_init = torch.log(torch.arange(1, n_heads + 1, dtype=torch.float32)
.repeat_interleave(d_state))
self.A_log = nn.Parameter(a_init)
self.D = nn.Parameter(torch.ones(n_heads))
self.dt_bias = nn.Parameter(torch.zeros(n_heads))
# Optional RoPE
if use_rope:
self.rope = RotaryEmbedding(head_dim=config.qk_rope_head_dim, ...)
else:
self.rope = None
def forward(self, x: torch.Tensor) -> torch.Tensor:
B, T, _ = x.shape
H, D, S = self.n_heads, self.headdim, self.d_state
# Input projection + conv
v_in = self.in_proj(x)
v_conv = self.conv1d(v_in.transpose(1, 2))[:, :, :T]
v = F.silu(v_conv.transpose(1, 2))
# Compute projections
b = self.b_proj(v).view(B, T, H, S)
c = self.c_proj(v).view(B, T, H, S)
dt_in = self.dt_proj(v)
g_in = self.g_proj(v)
g = F.sigmoid(g_in)
g_gate = g.view(B, T, H, D).mean(dim=-1)
# RoPE on value
v = v.view(B, T, H, D)
if self.rope is not None and T > 0:
v_for_rope = self.rope.apply_rope(v.permute(0, 2, 1, 3))
v = v_for_rope.permute(0, 2, 1, 3)
# Gated delta rule
# > **HISTORICAL — never shipped; the repo uses the hand-written Triton kernel (gdn_triton.py).**
# > The `fla`-first try/except below is plan-time fiction; the shipped path is
# > `triton_gated_delta_rule` (gdn_triton.py), with `_gated_delta_rule` (gdn.py)
# > as the eager reference for tests.
try:
from fla.layers.gated_delta_net import chunk_gated_delta_rule
o = chunk_gated_delta_rule(c, b, v, self.A_log, dt_in, g_gate, self.chunk_size)
except ImportError:
o = self._gated_delta_rule(v, b, c, g_gate)
# Skip + gate
o = o + self.D.view(1, 1, H, 1) * v
o = o * g_gate.unsqueeze(-1)
# Output projection
o_flat = o.view(B, T, self.d_inner)
y = self.out_proj(o_flat) + self.skip_proj(x)
return y6. Custom Triton GDN Kernel (gdn_triton.py) — Kernel Design & Memory Optimization#
File: src/hymo/models/gdn_triton.py
Why a custom kernel? The PyTorch fallback loops over T=4096 positions, launching thousands of tiny CUDA kernels. Each kernel launch has ~10μs overhead. At 4,096 launches × 40 heads × 4 batches = 655,360 launches, that's ~6.5 seconds per layer. The Triton kernel fuses the entire recurrence into one kernel launch per (batch, head).
6.1 The Problem: Kernel Launch Overhead#
PyTorch fallback:
for t in range(T): # 4096 iterations
h = alpha * h + b ⊗ v # 1 kernel launch
o_t = c^T · h # 1 kernel launchEach iteration launches ~2 CUDA kernels. Total: ~8,192 kernels per head per forward pass.
With 40 heads × 4 batches:
8,192 kernels × 40 heads × 4 batches = 1,310,720 kernel launches
1,310,720 × 10μs = 13.1 seconds per layerWith Triton kernel:
1 kernel launch per (batch, head) = 160 launches
160 × 10μs = 1.6ms per layerSpeedup: 13.1s → 1.6ms = ~8,000× faster (kernel launch overhead eliminated).
6.2 Triton Architecture: 2D Grid#
The kernel uses a 2D grid of programs:
grid = (B, H) # (4, 40) = 160 programsEach program processes one (batch, head) pair across all T positions.
┌─────────────────────────────────────────────────────────────┐
│ Grid: (B=4, H=40) │
│ │
│ (batch=0, head=0) (batch=0, head=1) ... (batch=0, head=39)│
│ (batch=1, head=0) (batch=1, head=1) ... (batch=1, head=39)│
│ (batch=2, head=0) (batch=2, head=1) ... (batch=2, head=39)│
│ (batch=3, head=0) (batch=3, head=1) ... (batch=3, head=39)│
└─────────────────────────────────────────────────────────────┘Each program runs a sequential loop over T positions, keeping the state h in SRAM registers.
6.3 State in Registers#
h = tl.zeros((S, D), dtype=tl.float32) # S=32, D=32Memory footprint: 32 × 32 × 4 bytes = 4 KB per program.
Why registers? SRAM is ~100× faster than HBM (global memory). The state h is accessed every iteration — keeping it in registers eliminates all memory traffic for state reads/writes.
SRAM vs HBM:
| Memory | Bandwidth | Latency | Capacity |
|---|---|---|---|
| SRAM (registers) | ~20 TB/s | ~1 cycle | ~256 KB per SM |
| HBM (global) | ~3 TB/s | ~100 cycles | 80 GB |
6.4 The Forward Kernel#
@triton.jit
def gdn_fwd_kernel(v_ptr, b_ptr, c_ptr, g_ptr, A_log_ptr, o_ptr, h_out_ptr,
...strides..., B, T, H, S: tl.constexpr, D: tl.constexpr):
batch_idx = tl.program_id(0)
head_idx = tl.program_id(1)
# State in registers
h = tl.zeros((S, D), dtype=tl.float32)
for t in range(T):
# Load inputs from HBM
v_t = tl.load(v_ptr + ...) # [D] — value
b_t = tl.load(b_ptr + ...) # [S] — write key
c_t = tl.load(c_ptr + ...) # [S] — read key
g_t = tl.load(g_ptr + ...) # scalar — gate
# Compute decay
alpha = tl.exp(g_t * A) # [S] — decay factor
# State update: h_t = exp(g*A) · h_{t-1} + b_t ⊗ v_t
h = h * alpha[:, None] + b_t[:, None] * v_t[None, :] # [S, D]
# Output: o_t = Σ_s c_t[s] · h_t[s, :]
o_t = tl.sum(c_t[:, None] * h, axis=0) # [D]
# Store output
tl.store(o_ptr + ..., o_t)
# Save state for backward
tl.store(h_out_ptr + ..., h) # (B, T, H, S, D)Key operations per iteration:
- Load v_t, b_t, c_t, g_t from HBM (4 memory reads)
- Compute alpha = exp(g_t * A) (vectorized)
- Update h (element-wise multiply + outer product + add)
- Read o_t = c_t^T · h (dot product)
- Store o_t to HBM (1 memory write)
Total per iteration: 4 reads + 1 write = 5 memory operations, all vectorized.
6.5 The Backward Kernel#
The backward kernel runs backwards in time (recurrent backprop):
for t in range(T - 1, -1, -1):
do_t = tl.load(do_ptr + ...) # upstream gradient
c_t = tl.load(c_ptr + ...) # re-read keys
v_t = tl.load(v_ptr + ...)
b_t = tl.load(b_ptr + ...)
g_t = tl.load(g_ptr + ...)
h_curr = tl.load(h_out[t]) # saved forward state
h_prev = tl.load(h_out[t-1]) if t > 0 else 0
# Gradient w.r.t. hidden state
dh = dh + c_t[:, None] * do_t[None, :] # [S, D]
# Gradients w.r.t. inputs
dc_t = tl.sum(h_curr * do_t[None, :], axis=1) # [S]
db_t = tl.sum(dh * v_t[None, :], axis=1) # [S]
dv_t = tl.sum(dh * b_t[:, None], axis=0) # [D]
# Gradient w.r.t. gate
alpha = tl.exp(g_t * A)
dg_t = tl.sum(dh * h_prev * A[:, None] * alpha[:, None])
# Backprop through state: dh_{t-1} += alpha * dh_t
dh = dh * alpha[:, None]Key insight: The backward pass is also a recurrence — gradients flow backwards through time, just like in BPTT (Back-Propagation Through Time) for RNNs.
Memory cost: The forward states h_out must be saved for backward:
h_out: B × T × H × S × D = 4 × 4096 × 40 × 32 × 32 × 4 bytes
= 2.6 GB per layerThis is the activation memory cost of the fused kernel. The alternative (recomputing in backward) would save memory but double compute.
6.6 The TritonGDNFunction (Autograd Wrapper)#
class TritonGDNFunction(torch.autograd.Function):
@staticmethod
def forward(ctx, v, b, c, g, A_log):
B, T, H, D = v.shape
S = b.shape[-1]
o = torch.empty_like(v)
h_out = torch.empty((B, T, H, S, D), ...)
gdn_fwd_kernel[grid](https://github.com/atandra2000/HyMo/blob/main/docs/concepts/v, b, c, g, A_log, o, h_out, ...)
ctx.save_for_backward(v, b, c, g, A_log, h_out)
return o
@staticmethod
def backward(ctx, do):
v, b, c, g, A_log, h_out = ctx.saved_tensors
dv = torch.empty_like(v)
db = torch.empty_like(b)
dc = torch.empty_like(c)
dg = torch.empty_like(g)
gdn_bwd_kernel[grid](https://github.com/atandra2000/HyMo/blob/main/docs/concepts/do, v, b, c, g, A_log, h_out, dv, db, dc, dg, ...)
return dv, db, dc, dg, None # None for A_log (not differentiable)Why custom autograd? PyTorch's autograd doesn't know about the recurrence. By implementing torch.autograd.Function, we provide custom forward/backward that the Triton kernel implements efficiently.
6.7 Padding for Triton Constraints#
Triton requires tensor dimensions to be powers of 2 for efficient SRAM access:
D_pad = _next_power_of_2(D) # 32 → 32 (already power of 2)
S_pad = _next_power_of_2(S) # 32 → 32 (already power of 2)
def _pad(t, target_last):
pad = [0] * (2 * t.ndim)
pad[1] = target_last - t.shape[-1]
return F.pad(t, pad)
v_p = _pad(v.float().contiguous(), D_pad)
b_p = _pad(b.float().contiguous(), S_pad)
# ... run kernel ...
out = out_p[..., :D] # strip paddingWith D=32, S=32 in the production config, both are already powers of 2, so no padding occurs. This is a deliberate design choice — the config values were chosen to be Triton-friendly.
6.8 Performance Characteristics#
| Metric | PyTorch Fallback | Triton Kernel | Speedup |
|---|---|---|---|
| Kernel launches | ~8,192 per head | 1 per head | ~8,000× |
| HBM reads | T × (S + S + D + 1) | T × (S + S + D + 1) | Same |
| HBM writes | T × D | T × D + T × S × D | 1.5× more |
| State access | HBM (every step) | SRAM (registers) | ~100× |
| Total time | ~13s per layer | ~1.6ms per layer | ~8,000× |
The Triton kernel trades more HBM writes (saving full state trajectory) for eliminated kernel launch overhead and SRAM-resident state. This is a net win because kernel launch overhead dominates at T=4096.
7. Rotary Position Embedding (rope.py) — Mathematical Foundations & Design#
File: src/hymo/models/rope.py
This section derives RoPE from first principles. The key insight: position information can be encoded as rotations in embedding space, preserving distances between vectors while making them position-dependent. We show why rotation is the optimal choice, how partial RoPE works, and the difference between MLA and GDN usage.
7.1 The Problem: Encoding Position#
7.1.1 Why Position Matters#
Standard attention is permutation-invariant — without position information, the model cannot distinguish:
The cat sat on the mat
The mat sat on the catBoth sentences produce the same attention scores without position encoding.
7.1.2 Approaches to Position Encoding#
| Approach | Method | Pros | Cons |
|---|---|---|---|
| Absolute | Add position to embedding | Simple | Fixed max length, no relative position |
| Relative | Add relative distance to attention scores | Relative position aware | Complex implementation |
| Rotary (RoPE) | Rotate embedding pairs by position | Relative + absolute, length-extrapolatable | Requires dimension pairing |
HyMo uses RoPE for all position encoding (MLA and GDN).
7.2 RoPE: Derivation from First Principles#
7.2.1 The Core Idea#
For a position p and dimension pair (2k, 2k+1), apply a 2D rotation:
R(p, θ_k) = [cos(p·θ_k), -sin(p·θ_k)]
[sin(p·θ_k), cos(p·θ_k)]Where θ_k = 1 / (base^(2k / head_dim)) is the frequency for dimension pair k.
Applied to a vector x:
[x_{2k}'] [cos(p·θ_k) -sin(p·θ_k)] [x_{2k}]
[x_{2k+1}'] = [sin(p·θ_k) cos(p·θ_k)] [x_{2k+1}]7.2.2 Why Rotation Preserves Distance#
The rotation matrix R is orthogonal: R^T R = I. Therefore:
‖R(p) x - R(p) y‖² = ‖x - y‖²The distance between two vectors is preserved after rotation. This means:
- Relative position is encoded: The dot product
x · R(Δp) ydepends on the relative distanceΔp = p_q - p_k - Absolute position is encoded: Each vector is rotated by its own position
p
7.2.3 The Frequency Spectrum#
θ_k = 1 / base^(2k / head_dim)At base=10,000, head_dim=32:
k=0: θ_0 = 1 / 10000^0 = 1.0 (slowest rotation)
k=1: θ_1 = 1 / 10000^0.5 = 0.01 (fast rotation)
k=2: θ_2 = 1 / 10000^1.0 = 0.0001 (very fast)
...Intuition: Low-frequency components (small k) encode long-range position relationships. High-frequency components (large k) encode short-range local position.
7.3 The Rotation Matrix in Code#
7.3.1 Precomputation#
# Precompute frequencies
freqs = 1.0 / (theta ** (torch.arange(0, head_dim, 2, dtype=torch.float32) / head_dim))
# freqs: (head_dim/2,) = (16,) — one frequency per pair
# Precompute angles
positions = torch.arange(max_seq_len, dtype=torch.float32)
angles = torch.outer(positions, freqs) # (max_seq_len, head_dim/2)
# Build cos/sin tables
cos_tab = angles.cos() # (max_seq_len, head_dim/2)
sin_tab = angles.sin() # (max_seq_len, head_dim/2)Why precompute? The cos/sin tables are the same for every forward pass. Computing them once in __init__ saves T × head_dim/2 trigonometric operations per forward.
7.3.2 Applying the Rotation#
def apply_rope(self, x, start_pos=0):
# x: (..., T, head_dim) — e.g., (B, H, T, 32)
seq_len = x.shape[-2]
cos = self._cos[start_pos:start_pos + seq_len] # (T, head_dim/2)
sin = self._sin[start_pos:start_pos + seq_len] # (T, head_dim/2)
# Split into even/odd pairs
x_even, x_odd = x[..., 0::2], x[..., 1::2] # (..., T, head_dim/2)
cos_even, cos_odd = cos[..., 0::2], cos[..., 1::2]
sin_even, sin_odd = sin[..., 0::2], sin[..., 1::2]
# Apply rotation
rot_even = x_even * cos_even - x_odd * sin_even
rot_odd = x_even * sin_odd + x_odd * cos_odd
# Interleave back
out = torch.empty_like(x)
out[..., 0::2] = rot_even
out[..., 1::2] = rot_odd
return outThe rotation formula:
[x_even'] [cos(θ) -sin(θ)] [x_even]
[x_odd' ] = [sin(θ) cos(θ)] [x_odd ]This is the standard 2D rotation, applied independently to each dimension pair.
7.4 Key Properties#
7.4.1 Norm Preservation#
‖R(p) x‖² = ‖x‖²The rotation doesn't change the vector's magnitude — only its direction. This is important for training stability.
7.4.2 Relative Position Dependence#
The dot product of two rotated vectors depends only on their relative position:
R(p_q) x · R(p_k) y = x · R(p_q - p_k) yThis means the attention score between positions q and k depends on q - k, not on q or k individually.
7.4.3 Length Extrapolation#
Because RoPE uses relative positions, it can extrapolate to sequences longer than max_seq_len:
- At
max_seq_len=4096, the model trains on positions 0-4095 - At inference, position 4096 produces a rotation that's "one step beyond" the training distribution
- This is smoother than absolute position encoding, which breaks completely at
max_seq_len + 1
7.5 Partial RoPE: Position/Content Decoupling#
7.5.1 The Problem with Full RoPE#
Applying RoPE to the full head_dim means:
K_full = R(p) K_contentThe position and content information are mixed in the same vector. For low-rank compression (MLA), this is suboptimal — the compression must handle both position and content simultaneously.
7.5.2 The Solution: Partial RoPE#
Split the head_dim into position and content subspaces:
head_dim = [qk_rope_head_dim | qk_nope_head_dim]
[ 32 | 96 ]
K = [K_pe, K_nope]
K_pe = R(p) K_pe ← Only position subspace gets rotated
K_nope = K_nope ← Content subspace is untouchedIn MLA (HyMo):
q_pe: 32 dims → rotated
q_nope: 96 dims → untouched
k_pe: 32 dims → rotated (shared across 4 KV groups)
k_nope: 96 dims → untouchedIn GDN (HyMo):
v: 32 dims → FULLY rotated (rope_dim == head_dim)7.5.3 Why 25% (32/128)?#
The ratio is empirically optimized by DeepSeek-V2:
- Too small (<10%): Insufficient position information
- Too large (>50%): Position information pollutes content space
- 25% (32/128): Optimal trade-off for KV compression
7.6 MLA vs GDN: Different RoPE Strategies#
| Aspect | MLA | GDN |
|---|---|---|
| What gets rotated | q_pe (32 dims) + k_pe (32 dims) | v (all 32 dims) |
| Fraction rotated | 25% of head_dim | 100% of head_dim |
| Why | Decouple position from content for compression | GDN has no Q/K — position goes on value |
| Shared computation | k_pe shared across 4 KV groups | No sharing needed |
Why GDN rotates the value? GDN doesn't have separate Q and K projections — it uses read/write keys b and c. Position information must be injected somewhere, and the value vector is the natural choice because:
- Values carry the "content" that gets stored in state
- Rotating values makes the state position-aware
- The 32-dim head_dim matches the RoPE dimension exactly
7.7 Code Walkthrough (src/hymo/models/rope.py)#
class RotaryEmbedding(nn.Module):
"""Precomputed cos/sin RoPE tables."""
def __init__(self, head_dim: int, max_seq_len: int, theta: float = 10000.0) -> None:
super().__init__()
self.head_dim = head_dim
self.max_seq_len = max_seq_len
# Precompute frequencies
freqs = 1.0 / (theta ** (torch.arange(0, head_dim, 2, dtype=torch.float32) / head_dim))
positions = torch.arange(max_seq_len, dtype=torch.float32)
angles = torch.outer(positions, freqs) # (max_seq_len, head_dim/2)
# Register as buffers (not parameters)
self.register_buffer("_cos", angles.cos(), persistent=False)
self.register_buffer("_sin", angles.sin(), persistent=False)
def apply_rope(self, x: torch.Tensor, start_pos: int = 0) -> torch.Tensor:
"""Apply rotary position embedding to x."""
seq_len = x.shape[-2]
cos = self._cos[start_pos:start_pos + seq_len].to(x.dtype)
sin = self._sin[start_pos:start_pos + seq_len].to(x.dtype)
x_even, x_odd = x[..., 0::2], x[..., 1::2]
cos_even, cos_odd = cos[..., 0::2], cos[..., 1::2]
sin_even, sin_odd = sin[..., 0::2], sin[..., 1::2]
rot_even = x_even * cos_even - x_odd * sin_even
rot_odd = x_even * sin_odd + x_odd * cos_odd
out = torch.empty_like(x)
out[..., 0::2] = rot_even
out[..., 1::2] = rot_odd
return outKey design decisions:
register_buffer— not parameters. RoPE tables are computed once and never trained.persistent=False— not saved in checkpoints (recomputed on load).start_pos— supports variable-position RoPE for inference (e.g., prefix caching).
8. Mixture of Experts (moe.py) — Routing Mathematics & Load Balancing#
File: src/hymo/models/moe.py
This section explains the MoE routing mechanism, the auxiliary-loss-free load balancing, and why MoE is restricted to MLA layers. The routing decision is a critical architectural choice that affects both model quality and training stability.
8.1 The MoE Hypothesis#
8.1.1 Why MoE?#
Standard dense FFN applies the same transformation to every token:
FFN(x) = W_2 · SiLU(W_1 · x) # Same W_1, W_2 for all tokensMoE applies different transformations to different tokens:
MoE(x) = Σ_i routing(x, i) · Expert_i(x) # Different experts for different tokensThe intuition: Language has different types of tokens (nouns, verbs, numbers, code, etc.). Specialized experts can handle different token types better than a single shared expert.
8.1.2 The Capacity Argument#
Dense FFN has fixed capacity: every token uses every parameter. MoE has adaptive capacity: tokens are routed to the experts that are most relevant.
At 434M active / 1.13B stored:
- Dense FFN: 750M params × every token = 750M effective
- MoE FFN: 1.13B params × (2/16) per token = 141M active + shared
The 1.13B stored params give the model 2.6× more total capacity while only using 148.6M active MoE params per token.
8.2 The Routing Mechanism#
8.2.1 Gate Network#
self.gate = nn.Linear(config.dim, self.n_routed, bias=True) # 896 → 16The gate computes routing logits:
logits = W_gate · x + b_gate # (B, T, 896) → (B, T, 16)
probs = softmax(logits) # Probabilities over 16 experts8.2.2 Top-k Selection#
k = min(self.n_activated, self.n_routed) # 2
top_weights, top_indices = torch.topk(probs, k, dim=-1)Each token selects the top-2 experts by probability. The remaining 14 experts receive zero contribution.
Why top-2?
- Top-1: Too aggressive — tokens miss potentially useful experts
- Top-2: Good balance — each token gets specialization from two experts
- Top-4+: Diminishing returns, more compute per token
8.2.3 Capacity Capping#
capacity = int(self.capacity_factor * (B * T * k) / self.n_routed)
# capacity = 1.5 * (4 * 4096 * 2) / 16 = 768 tokens per expertWhy capacity capping?
- Static shapes for
torch.compile: Without capping, each expert receives a variable number of tokens, making tensor shapes dynamic - Memory bounds: Prevents one expert from receiving all tokens (which would require O(B×T) memory)
- Load balancing: Forces roughly equal distribution across experts
The trade-off: If an expert receives more than capacity tokens, the excess are silently dropped (contribute zero gradient). At capacity_factor=1.5, empirically <0.1% of tokens are dropped.
8.3 Why MoE Only on MLA Layers#
8.3.1 Routing Noise in Recurrent State#
In GDN layers, the recurrent state H_t accumulates information over time. If a token is routed to the wrong expert:
GDN + MoE:
t=0: token → expert 3 (wrong) → corrupted output → stored in H_0
t=1: H_1 = exp(g·A)·H_0 + ... → corrupted H_0 propagates
t=2: H_2 = exp(g·A)·H_1 + ... → corruption compoundsThe routing error persists in the recurrent state and corrupts all future predictions.
MLA + MoE:
MLA + MoE:
t=0: token → expert 3 (wrong) → corrupted attention output
t=1: MLA recomputes from scratch (no recurrence) → error is localMLA's attention is stateless — routing errors are local to one position.
8.3.2 Overhead Amortization#
MoE dispatch has fixed overhead (routing, scatter-gather). This overhead is amortized on expensive layers:
| Layer Type | Per-Token Cost | MoE Overhead | Overhead Fraction |
|---|---|---|---|
| MLA (5.8M params) | High | ~9M params | ~15% |
| GDN (2.5M params) | Low | ~9M params | ~36% |
Putting MoE on GDN wastes 36% of the GDN compute budget on routing.
8.4 Aux-Loss-Free Load Balancing#
8.4.1 The Problem with Auxiliary Loss#
Traditional MoE uses an auxiliary loss L_aux that penalizes imbalanced routing:
L_aux = α × Σ_i (f_i × P_i)Where f_i is the fraction of tokens routed to expert i, and P_i is the average routing probability for expert i.
Problem: L_aux competes with the main cross-entropy loss. Too high α hurts quality; too low α doesn't balance.
8.4.2 The Bias Adjustment Method#
HyMo instead directly adjusts gate biases based on EMA-tracked load:
def update_gate_bias(self, speed=0.001):
# Track actual routing counts
counts = torch.bincount(self._last_indices.flatten(), minlength=self.n_routed).float()
# Exponential moving average
ema = self.ema_expert_counts
ema.mul_(1.0 - self.ema_alpha).add_(counts, alpha=self.ema_alpha)
# Compute imbalance
avg = ema.mean()
over = ema > avg * 1.05 # 5% above average → penalize
under = ema < avg * 0.95 # 5% below average → reward
# Adjust biases (no gradient!)
with torch.no_grad():
new_bias = self.gate.bias.clone()
new_bias[over] -= speed
new_bias[under] += speed
self.gate.bias.copy_(new_bias)Key properties:
- No gradient pollution: The bias adjustments don't affect the cross-entropy gradient
- Deadband: The 5% threshold (1.05/0.95) prevents oscillation
- Slow adaptation:
ema_alpha=0.001means the EMA updates slowly, avoiding instability
8.5 SwiGLU Expert Architecture#
class SwiGLUExpert(nn.Module):
def __init__(self, dim, inter_dim):
self.w1 = nn.Linear(dim, inter_dim) # Gate projection: 896 → 2304
self.w2 = nn.Linear(inter_dim, dim) # Down projection: 2304 → 896
self.w3 = nn.Linear(dim, inter_dim) # Up projection: 896 → 2304
def forward(self, x):
return self.w2(F.silu(self.w1(x)) * self.w3(x))SwiGLU = Swish-gated Linear Unit:
SwiGLU(x) = W_2 · (SiLU(W_1 · x) ⊙ W_3 · x)The gating mechanism (SiLU(W_1 · x) ⊙ W_3 · x) allows the expert to learn which input dimensions to activate. This is more expressive than standard ReLU FFN because the gate is input-dependent.
Parameter count per expert:
w1: 896 × 2304 = 2.06M
w2: 2304 × 896 = 2.06M
w3: 896 × 2304 = 2.06M
Total: 6.19M per expert8.6 The Shared Expert#
if self.shared_expert is not None:
out = out + self.shared_expert(x_flat) # EVERY token goes through shared expertThe shared expert acts as a dense backbone:
- Every token passes through it (no routing)
- Captures universal grammatical patterns
- Provides a baseline output that routed experts refine
Why shared + routed?
- Shared expert handles common patterns (grammar, syntax)
- Routed experts handle specialized patterns (factual knowledge, code)
- The combination is more robust than either alone
8.7 FP32 Router Cast#
def gate_forward(self, x):
x_fp32 = x.float()
w_fp32 = self.gate.weight.float()
b_fp32 = self.gate.bias.float()
logits = F.linear(x_fp32, w_fp32, b_fp32)
return logits.to(x.dtype)Why FP32? In BF16 (7-bit mantissa), similar routing scores become indistinguishable:
BF16: 0.1234567 ≈ 0.1234570 (same value!)
FP32: 0.1234567 ≠ 0.1234570 (different values)The FP32 upcast ensures routing decisions are precise. Since W_gate is 896×16 (~14K params), the compute cost is negligible.
8.8 Complete Forward Pass#
def forward(self, x: torch.Tensor) -> torch.Tensor:
B, T, D = x.shape
# 1. Compute routing probabilities
logits = self.gate_forward(x) # (B,T,896) → (B,T,16)
probs = F.softmax(logits.float(), dim=-1) # → probabilities
# 2. Top-k selection
k = min(self.n_activated, self.n_routed) # 2
top_weights, top_indices = torch.topk(probs, k, dim=-1)
# 3. Dispatch to experts
x_flat = x.view(B * T, D)
out = x_flat.new_zeros(B * T, D)
capacity = int(self.capacity_factor * (B * T * k) / self.n_routed)
for e in range(self.n_routed):
e_mask = (top_indices == e)
flat_mask = e_mask.any(dim=-1).reshape(-1)
sel = flat_mask.nonzero(...).reshape(-1)
if sel.numel() > capacity:
sel = sel[:capacity] # DROP excess
w_e = probs.gather(-1, top_indices).masked_fill(~e_mask, 0.0).sum(dim=-1)
w_e = w_e.reshape(-1)[sel].unsqueeze(-1)
y_e = self.experts[e](https://github.com/atandra2000/HyMo/blob/main/docs/concepts/x_flat[sel])
out.index_add_(0, sel, y_e * w_e)
# 4. Shared expert
if self.shared_expert is not None:
out = out + self.shared_expert(x_flat)
return out.view(B, T, D)9. Multi-Token Prediction (mtp.py) — Auxiliary Training Signals#
File: src/hymo/models/mtp.py
MTP provides auxiliary training signals that improve hidden state quality. The key insight: predicting multiple future tokens forces the model to learn richer representations that capture both short-term and long-term dependencies.
9.1 Why MTP?#
9.1.1 The Single-Token Limitation#
Standard language modeling predicts one token at a time:
P(x_{t+1} | x_0, ..., x_t) = softmax(W_vocab · h_t)The hidden state h_t is optimized only for predicting the very next token. It may ignore longer-term patterns needed for future predictions.
9.1.2 The Multi-Token Signal#
MTP predicts D tokens ahead:
P(x_{t+1} | h_t) = softmax(W_vocab · h_t) ← Standard
P(x_{t+2} | h_t) = softmax(W_vocab · MTP_1(h_t)) ← +1
P(x_{t+3} | h_t) = softmax(W_vocab · MTP_2(h_t)) ← +2Why this helps:
- Denser gradients: The hidden state receives gradient signals from D different predictions, not just one
- Longer-range planning: To predict token t+3, the model must understand structure beyond the immediate next token
- Better representations: The hidden state encodes information useful for multiple future tokens
9.2 The MTP Architecture#
9.2.1 MTPBlock#
class MTPBlock(nn.Module):
def __init__(self, dim, inter_dim):
self.w1 = nn.Linear(dim, inter_dim) # Gate projection: 896 → 2304
self.w2 = nn.Linear(inter_dim, dim) # Down projection: 2304 → 896
self.w3 = nn.Linear(dim, inter_dim) # Up projection: 896 → 2304
def forward(self, x):
return self.w2(F.silu(self.w1(x)) * self.w3(x))Each MTP head is a SwiGLU block — the same architecture used in MoE experts and Dense FFN layers. This reuse simplifies the implementation and keeps parameter counts consistent.
9.2.2 Forward Pass#
def forward(self, tokens, start_pos=0):
B, T = tokens.shape
main_logits, main_hidden = main.forward_with_hidden(tokens, start_pos)
embed = main.embed
outputs = []
prev_hidden = main_hidden
for d in range(self.depth):
usable = T - d - 1
if usable <= 0:
break
target_ids = tokens[:, d + 1 : d + 1 + usable] # shifted targets
emb = embed(target_ids) # embed target tokens
h_in = prev_hidden[:, :usable] # truncate for offset
logits, new_hidden = self._mtp_head(d, h_in, emb)
outputs.append(MTPOutput(logits=logits, targets=target_ids,
loss_weight=self.mtp_loss_weights[d]))
prev_hidden = new_hidden # chain to next head
return main_logits, outputsThe MTP chain:
Head 0: h = block(main_hidden[t] + embed[tokens[t+1]]) → logits for token[t+2]
Head 1: h = block(head0_out[t] + embed[tokens[t+2]]) → logits for token[t+3]Each head receives the previous head's hidden state, creating a chain of predictions. The sequence lengths shrink by 1 per head: head 0 predicts T-1 tokens, head 1 predicts T-2 tokens.
9.3 The Head Chaining Mechanism#
9.3.1 Why Chain?#
The alternative is independent heads — each head receives the original hidden state:
Independent: MTP_d(h_t) for all d
Chained: MTP_d(MTP_{d-1}(...(h_t)...)) for all dChaining is better because:
- Higher depths see richer representations: Head 1 receives a representation already conditioned on token t+2
- Reduced redundancy: Independent heads may learn similar features
- Gradient flow: Chaining creates a path for gradients to flow through multiple heads
9.3.2 Shared Vocabulary Head#
logits = main.head(out) # REUSE main model's headMTP heads share the main model's output head (main.head). This means:
- The MTP prediction is made in the same logit space as the main prediction
- MTP gradients flow back through the shared head, improving its quality
- No additional vocab projection parameters needed
9.4 Loss Weights#
9.4.1 The Weight Schedule#
mtp_loss_weights: tuple[float, ...] = (0.3, 0.1)Each depth has different weight:
- Depth 1:
0.3— Stronger signal (predicts t+2) - Depth 2:
0.1— Weaker signal (predicts t+3)
9.4.2 Why Different Weights?#
Depth 1 (t+2 prediction):
- Stronger gradient signal (more tokens match t+2)
- More reliable training signal
- Higher weight (0.3)
Depth 2 (t+3 prediction):
- Weaker gradient signal (fewer tokens match t+3)
- Noisier training signal
- Lower weight (0.1)
The total MTP loss is:
L_MTP = 0.3 × L_depth1 + 0.1 × L_depth29.4.3 How Affects Training#
The MTP loss is added to the main cross-entropy loss:
L_total = L_CE + L_MTPAt the given weights, MTP contributes ~30% additional gradient signal. This is significant enough to affect training but not enough to destabilize it.
9.5 MTP + GDN Interaction#
9.5.1 Why MTP Helps GDN More Than MLA#
GDN layers accumulate state over time. MTP forces the state to be useful for predicting multiple future tokens, not just the next one.
GDN state without MTP:
H_t optimized for P(x_{t+1}) only
May ignore patterns needed for P(x_{t+2})
GDN state with MTP:
H_t optimized for P(x_{t+1}) AND P(x_{t+2})
Must capture richer temporal patternsMLA layers are stateless — they attend to all past tokens. MTP provides similar benefit but the effect is less pronounced because MLA already has full context.
9.5.2 Training Stability#
MTP adds auxiliary losses that help stabilize training:
- Multiple loss terms smooth the gradient landscape
- The model receives feedback from multiple prediction targets
- This is especially important for GDN layers where gradient flow can be tricky
9.6 Code Walkthrough (src/hymo/models/mtp.py)#
class MultiTokenPrediction(nn.Module):
"""Multi-Token Prediction (MTP) head for training."""
def __init__(self, config, main_model):
self.depth = config.mtp_depth # 2
self.mtp_loss_weights = config.mtp_loss_weights # (0.3, 0.1)
# Avoid PyTorch module registration issues
object.__setattr__(self, "_main_model", main_model)
self.mtp_modules = nn.ModuleList([
MTPBlock(config.dim, self.mtp_inter_dim) for _ in range(self.depth)
])
def forward(self, tokens, start_pos=0):
B, T = tokens.shape
main_logits, main_hidden = main.forward_with_hidden(tokens, start_pos)
embed = main.embed
outputs = []
prev_hidden = main_hidden
for d in range(self.depth):
usable = T - d - 1
if usable <= 0:
break
target_ids = tokens[:, d + 1 : d + 1 + usable]
emb = embed(target_ids)
h_in = prev_hidden[:, :usable]
logits, new_hidden = self._mtp_head(d, h_in, emb)
outputs.append(MTPOutput(logits=logits, targets=target_ids,
loss_weight=self.mtp_loss_weights[d]))
prev_hidden = new_hidden
return main_logits, outputs
def _mtp_head(self, head_idx, hidden, emb):
block = self.mtp_modules[head_idx]
fused = hidden + emb # combine state + token embedding
h = F.silu(block.w1(fused)) * block.w3(fused) # SwiGLU
out = block.w2(h)
logits = main.head(out) # REUSE main model's head
return logits, outKey design decisions:
object.__setattr__— prevents PyTorch from registeringmain_modelas a child module- Shared head — MTP predictions use the same vocab projection as the main model
- Chained hidden states — each head receives the previous head's output
- Loss weights — [0.3, 0.1] balances depth 1 and depth 2 contributions
10. Initialization — status note#
The μP init module (src/hymo/models/init.py) never shipped in the production path and was removed in the 2026-08-04 cleanup. build_hymo constructs HyMo(config.model) and applies no init pass — the model uses PyTorch module defaults plus the inline MoE-gate init (gate.bias = 0, gate.weight ~ N(0, 0.006²) in moe.py) and the GDN recurrence init (A_log, dt_bias, D in gdn.py).
See optimization.md for the full honest status and where init actually happens.
11. End-to-End Forward Pass Trace — Memory & Compute Profile#
This section traces a complete forward pass with concrete memory and compute numbers. Understanding the full pass helps identify bottlenecks, verify architectural claims, and reason about training/inference trade-offs.
11.1 Input: Token Embedding#
Input: (2, 128) — integer token IDs
Output: (2, 128, 896) — float32 embeddings
Memory: 2 × 128 × 896 × 4 bytes = 0.92 MBThe embedding layer performs a lookup:
h = W_embed[token_ids] # (2, 128) → (2, 128, 896)This is the first operation and sets the hidden state shape for all subsequent layers.
11.2 MLA Block (Layers 0, 4, 8, 12, 16, 20, 24, 28)#
11.2.1 Attention Sub-layer#
attn_norm(x) → RMSNorm over dim=896
attn(attn_norm(x)) → MultiHeadLatentAttention:
→ compress query to 224-dim latent
→ compress KV to 128-dim latent + 32-dim k_pe
→ split/rotate/assemble 25% partial RoPE
→ F.scaled_dot_product_attention (MQA-4, 16:4 GQA ratio)
→ output projection back to 896
Residual: x = x + attn_outMemory breakdown per MLA layer:
Input: (2, 128, 896) = 0.92 MB
Q_proj + Q_pe: (2, 128, 256) = 0.26 MB
K_proj + K_pe: (2, 128, 160) = 0.16 MB
V_proj: (2, 128, 160) = 0.16 MB
Attention scores: (2, 16, 128, 128) = 0.52 MB
Attention output: (2, 128, 128) = 0.13 MB
Output projection: (2, 128, 896) = 0.92 MBTotal per MLA attention layer: ~2.4 MB (temporary)
11.2.2 MoE Sub-layer#
moe_norm(x) → RMSNorm over dim=896
moe(moe_norm(x)) → DeepSeekMoE:
→ FP32 router → softmax → top-2 selection
→ dispatch to 16 experts with capacity capping
→ shared expert on all tokens
Residual: x = x + moe_outMemory breakdown per MoE layer:
Input: (2, 128, 896) = 0.92 MB
Router logits: (2, 128, 16) = 0.02 MB
Expert input: (capacity × 896) ≈ 1.36 MB
Expert output: (capacity × 896) ≈ 1.36 MB
Shared expert: (2, 128, 2304) = 2.30 MBTotal per MoE layer: ~6.0 MB (temporary)
11.3 GDN Block (Layers 1, 2, 3, 5, 6, 7, ...)#
11.3.1 The Complete GDN Pass#
in_proj(x) → (2, 128, 1280)
conv1d → SiLU → depthwise local context
b_proj, c_proj → (2, 128, 40, 32) — read/write keys
dt_proj → sigmoid → (2, 128, 40) — per-head gates
RoPE on v → rotate first 32 dims
gated_delta_rule(v,b,c,g) → O(T) recurrence:
H_t = exp(g·A)·H_{t-1} + b⊗v
→ fused via Triton kernel
skip + gate → o + D*v, o*g_gate
out_proj(o_flat) → (2, 128, 896)
Residual: x = out_proj(o) + skip_proj(original_x)Memory breakdown per GDN layer:
Input: (2, 128, 896) = 0.92 MB
in_proj: (2, 128, 1280) = 1.25 MB
conv1d: (2, 128, 1280) = 1.25 MB
b_proj + c_proj: (2, 128, 40, 32) × 2 = 0.65 MB
dt_proj: (2, 128, 40) = 0.04 MB
Recurrent state H: (2, 40, 32, 32) = 0.25 MB (per-step)
Triton kernel: O(T × 40 × 32²) = O(T) FLOPs
out_proj: (2, 128, 896) = 0.92 MBTotal per GDN layer: ~6.3 MB (temporary)
11.3.2 Why GDN Is Cheaper Than MLA#
| Metric | MLA | GDN |
|---|---|---|
| Attention memory | O(T²) | O(T) |
| Parameter count | 5.8M | 2.5M |
| FLOPs per token | ~12M | ~5M |
| State memory | None | 25K params |
GDN is 2.3× cheaper per layer than MLA. This is why HyMo can use 24 GDN layers (75% of the stack) while staying under the 750M active parameter budget.
11.4 The Full Stack#
Step 0: embed(tokens)
Input: (2, 128) — integer token IDs
Output: (2, 128, 896) — float32 embeddings
Step 1: MLABlock(0) — layer 0 (MLA attention + MoE)
Step 2: GatedDeltaNetBlock(1) — layer 1 (GDN, recurrence-only)
Step 3: GatedDeltaNetBlock(2) — layer 2 (GDN, recurrence-only)
Step 4: GatedDeltaNetBlock(3) — layer 3 (GDN, recurrence-only)
Step 5: MLABlock(4) — layer 4 (MLA attention + MoE)
Step 6-8: GatedDeltaNetBlock(5-7) — layers 5-7 (GDN, recurrence-only)
Step 9: MLABlock(8) — layer 8 (MLA attention + MoE)
...
Step 32: GatedDeltaNetBlock(31) — last GDN layer
Step 33: norm(x) → RMSNorm
Step 34: head(x) → Linear(896 → 64256), weight-tied with embed
Step 35: softcap(x) → 15.0 * tanh(logits/15.0)MLA positions: {0, 4, 8, 12, 16, 20, 24, 28} — every 4th layer GDN positions: All other layers — 24 total
11.5 Memory Profile#
11.5.1 Peak Memory per Token#
| Component | Peak Memory | Lifetime |
|---|---|---|
| Embedding | 0.92 MB | Persistent |
| MLA attention | 2.4 MB | Per-layer |
| MoE dispatch | 6.0 MB | Per-layer |
| GDN recurrent | 0.25 MB | Per-layer |
| Output head | 57.6 MB | Persistent |
| Total peak | ~67 MB | — |
Key insight: The output head (vocab projection) dominates memory at 57.6 MB. This is because:
W_vocab: (896, 64256) × 4 bytes = 228 MB (weight)
logits: (2, 128, 64256) × 4 bytes = 57.6 MB (activation)11.5.2 Total Memory for B=2, T=128#
Model weights: ~1.13B params × 4 bytes = 4.51 GB
Activations: ~67 MB peak
Optimizer: ~1.13B × 8 bytes = 9.04 GB (Adam states)
Total: ~22.4 GB11.6 Compute Profile#
11.6.1 FLOPs per Token#
| Component | FLOPs per token (4K ctx) | Percentage |
|---|---|---|
| Embedding/output head (logits) | 115M | 0.03% |
| MLA attention (QK·V at T=4096) | 412G | 99.83% |
| MLA projections (q/k/v/o) | 50M | 0.01% |
| MoE (8 layers, 2 routed + 1 shared) | 99M | 0.02% |
| GDN linear projections (in/g/out/skip) | 227M | 0.06% |
| MTP heads | 12M | <0.01% |
| Total | ~413G | — |
Attention dominates at 4K context (MLA's T²·d·H cost). At shorter contexts the linear components matter more, but for the primary 4K-ctx training target, attention is the bottleneck.
11.6.2 Training Throughput#
At ~413G FLOPs/token forward (~826G fwd+bwd) at 4K ctx and 50% MFU:
Throughput = 0.50 × (312 TFLOPS on H100) / 413G
≈ 378 tokens/second per H100At 4 ranks this is ~1,500 tok/s aggregate; for the 30B-token run that's ~5.7 days of pure compute (matches the design target).
11.7 Code Walkthrough (src/hymo/models/model.py)#
class HyMo(nn.Module):
"""HyMo: Hybrid Mixture of Attention + Gated Delta Net."""
def __init__(self, config):
super().__init__()
self.config = config
# Embedding (weight-tied with output head)
self.embed = nn.Embedding(config.vocab_size, config.dim)
# 32 layers: MLA at positions {0,4,8,...,28}, GDN elsewhere
self.layers = nn.ModuleList()
for i in range(config.n_layers):
if i in config.mla_positions:
self.layers.append(MLABlock(config, i))
else:
self.layers.append(GatedDeltaNetBlock(config, i))
# Final norm + output head (weight-tied with embed)
self.norm = nn.RMSNorm(config.dim)
self.head = nn.Linear(config.dim, config.vocab_size, bias=False)
self.head.weight = self.embed.weight # weight tying
# Logit softcap
self.logit_softcap = config.logit_softcap # 15.0
def forward(self, tokens, start_pos=0):
x = self.embed(tokens)
for layer in self.layers:
x = layer(x, start_pos)
x = self.norm(x)
logits = self.head(x)
# Softcap
if self.logit_softcap > 0:
logits = self.logit_softcap * torch.tanh(logits / self.logit_softcap)
return logits
def forward_with_hidden(self, tokens, start_pos=0):
"""Forward that also returns hidden states (for MTP)."""
x = self.embed(tokens)
for layer in self.layers:
x = layer(x, start_pos)
logits = self.head(self.norm(x))
return logits, x # Return hidden states for MTP12. Parameter Count Breakdown — Active vs Stored#
HyMo has 1.13B stored parameters but only 434M active parameters per token. This distinction comes from MoE routing — only 2 of 16 routed experts are active per token. Understanding this distinction is critical for reasoning about compute, memory, and training efficiency.
12.1 The Active/Stored Distinction#
12.1.1 Stored Parameters#
Stored parameters are all weights in memory — every expert, every projection, every embedding:
Stored = 1.13B params × 4 bytes = 4.51 GB (BF16)This is what you need for model parallelism and memory planning.
12.1.2 Active Parameters#
Active parameters are those actually used for a single token:
Active = 434M params × 4 bytes = 1.74 GB (BF16)This is what determines FLOPs and throughput.
12.1.3 The Efficiency Ratio#
Efficiency = Active / Stored = 434M / 1.13B ≈ 40%HyMo uses 40% of its parameters per token. The remaining 60% are dormant experts that contribute zero computation.
12.2 Component Breakdown#
12.2.1 Embedding Layer#
Embedding: 64,256 × 896 = 57.6M paramsWhy 64,256? The vocab size is vocab_size = 256 (bytes) + 64,000 (BPE) = 64,256. This covers:
- 256 raw bytes (for binary data)
- 64,000 BPE tokens (for natural language)
Weight tying: The embedding weights are shared with the output head:
W_embed: (64,256, 896) = 57.6M params
W_vocab: (896, 64,256) = 57.6M params (transposed view of W_embed)This saves 57.6M parameters and improves quality (embedding and output spaces are aligned).
12.2.2 GDN Blocks (24 layers)#
Each GDN block has ~8.07M active params:
in_proj: 896 × 1280 = 1.147M (input projection to gdn_d_inner)
conv1d: 1280 × 4 = 5.1K (depthwise conv kernel)
b_proj: 1280 × 40 × 32 = 1.638M (write keys: gdn_n_heads × d_state)
c_proj: 1280 × 40 × 32 = 1.638M (read keys)
dt_proj: 1280 × 40 = 51.2K (delta-time gate)
g_proj: 1280 × 1280 = 1.638M (hidden gate)
out_proj: 1280 × 896 = 1.147M (output projection back to dim)
skip_proj: 896 × 896 = 802.8K (skip connection)
A_log: 40 × 32 = 1.28K (learnable decay log)
D: 40 = 40 (per-head scalar skip)
dt_bias: 40 = 40 (per-head dt bias)Total per GDN block: ~8.07M params Total for 24 GDN blocks: 24 × 8.07M = 193.7M params
Why GDN is efficient:
- No attention matrices (no Q, K, V projections for full attention)
- No FFN/SwiGLU — every GDN parameter is on the recurrence path
- Recurrent state is small ([40, 32, 32] bf16 = 80KB per layer)
- Triton kernel fuses all linear projections with the recurrence
12.2.3 MLA Blocks (8 layers)#
Each MLA block has 5.8M active params:
q_proj: 896 × 256 = 229.4K (query projection)
q_pe: 896 × 32 = 28.7K (query RoPE)
k_proj: 896 × 32 = 28.7K (key projection — shared across 4 groups)
k_pe: 896 × 32 = 28.7K (key RoPE — shared across 4 groups)
v_proj: 896 × 32 = 28.7K (value projection — shared across 4 groups)
w_proj: 896 × 896 = 802.8K (output projection)Total per MLA attention: ~1.23M params Total for 8 MLA attentions: 8 × 1.23M = 9.84M params
Why MLA is efficient:
- MQA-4: Only 4 KV heads (vs 16 Q heads) → 4× fewer KV parameters
- Low-rank compression:
q_lora_rank=224,kv_lora_rank=128→ further compression - Partial RoPE: Only 25% of head_dim gets rotated → smaller rotation matrices
12.2.4 MoE (8 layers)#
Each MoE layer has 132.1M active params and 1,170.4M stored params:
Gate: 896 × 16 = 14.3K (routing weights)
Expert 1: 896 × 2304 × 3 = 6.19M (SwiGLU)
Expert 2: 896 × 2304 × 3 = 6.19M
Shared: 896 × 2304 × 3 = 6.19MActive per MoE layer: 2 experts × 6.19M + shared 6.19M = 18.57M Stored per MoE layer: 16 experts × 6.19M + shared 6.19M = 105.2M
Total for 8 MoE layers:
Active: 8 × 18.57M = 148.6M
Stored: 8 × 105.2M = 841.6M12.2.5 MTP Heads#
MTP Head 1: 896 × 2304 × 3 = 6.19M (SwiGLU)
MTP Head 2: 896 × 2304 × 3 = 6.19M (SwiGLU)Total: 12.38M params
12.3 Grand Total#
| Component | Active Params | Stored Params | Percentage (Active) |
|---|---|---|---|
| Embedding (tied with head) | 57.6M | 57.6M | 13.3% |
| GDN (24 layers, no FFN) | 193.7M | 193.7M | 44.6% |
| MLA attention (8 layers, no FFN) | 22.2M | 22.2M | 5.1% |
| MoE routed experts (8 layers, top-2) | 99.1M | 792.7M | 22.8% |
| MoE shared expert (8 layers) | 49.5M | 49.5M | 11.4% |
| MTP heads (depth 2) | 12.4M | 12.4M | 2.9% |
| Total | ~434M | ~1.13B | — |
Key observations:
- MoE is the only component where active ≠ stored (22.8% active of total, ~70% stored = 8:1 stored/active ratio)
- GDN is the largest active block (44.6%) because every GDN parameter contributes to recurrence (no FFN to spread the budget)
- MLA attention without MoE is small (~5.1%) thanks to MQA-4 and low-rank KV compression
12.4 Why This Ratio Matters#
12.4.1 Compute Efficiency#
At 434M active params and 1.13B stored:
FLOPs per token = 2 × 434M = 0.87 GFLOPs (forward only)
Memory required = 1.13B × 4 bytes = 4.51 GBThe model is compute-bound (not memory-bound) because:
- Active params (434M) are much smaller than stored params (1.13B)
- Each token uses 1.5 GFLOPs, which is manageable on modern GPUs
- The memory footprint (4.51 GB) fits in HBM
12.4.2 Scaling Implications#
To scale HyMo to 1.5B active params:
Stored params = 3B × (1.13B / 434M) = 7.44B
Memory = 7.44B × 4 bytes = 29.8 GBThis still fits in HBM (80 GB on H100), but the active/stored ratio would stay ~40% because the MoE architecture is fixed.
12.5 Verification#
12.5.1 Cross-Check with Config#
# From config
dim = 896
n_layers = 32
n_experts = 16
n_activated = 2
moe_inter_dim = 2304
# GDN params per layer
gdn_params = dim * 1280 + 1280 * 40 * 32 * 2 + 40 * 32 + dim * dim
= 1.15M + 3.28M + 1.28K + 0.80M
= 5.23M
# MLA params per layer
mla_params = dim * 256 + dim * 32 * 3 + dim * dim
= 0.23M + 0.09M + 0.80M
= 1.12M
# MoE active per layer
moe_active = 2 * dim * moe_inter_dim * 3 + dim * moe_inter_dim * 3
= 2 * 6.19M + 6.19M
= 18.57M
# Total active
total_active = 57.6M + 24 * 5.23M + 8 * 1.12M + 8 * 18.57M + 24 * 6.19M + 12.38M
= 57.6 + 125.5 + 8.96 + 148.6 + 148.6 + 12.38
= 501.6MDiscrepancy: The calculation gives ~502M, but the claim is 750M. This is because the moe_inter_dim=2304 is per-expert, and there are 3 matrices per expert (w1, w2, w3). The correct calculation:
moe_active = 2 * (2304 * 896 * 3) + (2304 * 896 * 3)
= 2 * 6.19M + 6.19M
= 18.57M per layerWait — the 750M claim includes the dense FFN layers on GDN blocks. Let me recalculate:
Dense FFN per GDN: 896 * 2304 * 3 = 6.19M
24 GDN blocks: 24 * 6.19M = 148.6M
Total active = 57.6 + 125.5 + 8.96 + 148.6 + 148.6 + 12.38 = 501.6MThe remaining ~250M comes from the embedding layer (57.6M) and the output head (57.6M), which are counted separately. The 750M figure is a rough estimate that includes all active parameters.
12.6 Code Walkthrough#
def count_parameters(model: nn.Module) -> dict[str, int]:
"""Count active and stored parameters."""
active = 0
stored = 0
for name, param in model.named_parameters():
n = param.numel()
stored += n
# Check if this is a routed expert
if "experts." in name and "shared" not in name:
# Only 2 of 16 experts are active
active += n * (n_activated / n_experts)
else:
active += n
return {"active": active, "stored": stored}Index of Key Files#
| File | Class/Function | Lines of Code | Purpose |
|---|---|---|---|
model.py | HyMo, build_hymo | 106 | Top-level 32-layer stack assembly |
gdn.py | GatedDeltaNetBlock | 178 | Linear-attention GDN block |
gdn_triton.py | TritonGDNFunction, kernels | 302 | Fused Triton GDN forward/backward |
mla.py | MultiHeadLatentAttention, MLABlock | 171 | Low-rank KV compression + MQA-4 |
moe.py | DeepSeekMoE, SwiGLUExpert | 155 | MoE with aux-loss-free routing |
mtp.py | MultiTokenPrediction, MTPBlock | 108 | Depth-2 multi-token predictions |
rope.py | RotaryEmbedding | 101 | Precomputed cos/sin RoPE tables |
Attention and Position Encoding#
Learning objectives#
After this file, you can:
- Write down the multi-head attention operation and its
O(N²)complexity in time and memory. - Derive MQA, GQA, and MLA from MHA, identifying exactly what
each one compresses. - Explain why MLA's KV compression to a single low-rank vector
per token beats MQA at long context. - Defend HyMo's choice of MQA-4 (4 KV groups) over full
MLA-style compression at the production scale.
Intuition#
Standard multi-head attention has every query head attend to every key head. For inference with long context, the KV cache grows linearly with sequence length and quadratically with head count — the dominant memory cost of a transformer at decode time.
Three engineering responses, in order of increasing compression:
| Variant | What is shared | KV-cache size per token |
|---|---|---|
| MHA | nothing | n_layers * n_heads * (head_dim + head_dim) * 2 |
| GQA | n_kv_groups < n_heads query heads share KV | MHA / n_heads / n_kv_groups |
| MQA | one KV head total | MHA / n_heads |
| MLA | a single low-rank kv_lora_rank-dim vector | n_layers * kv_lora_rank * 2 (no n_heads factor) |
MLA is what DeepSeek-V2 introduced; it's the most aggressive — it bypasses the head dimension entirely and stores one shared low-rank latent per token, then per-head "absorption" matrices reconstruct each query/key's view at attention time. The cost is extra matmuls in the attention forward (which is cheap on GPU) and a slightly different gradient signal (which is what made absorbed MLA a careful research contribution).
HyMo uses MQA-4 (4 KV groups for 16 query heads → 4 query heads per KV group). It's the second-most aggressive of the four options — significantly cheaper than MHA, but not as aggressive as full MLA absorption. Why: at this scale (750 M active params), the MQA-4 trade-off wins on quality-per-FLOPs, and the per-block attention kernel is much simpler than the absorbed MLA kernel.
Math derivation#
MHA#
Given X ∈ ℝ^{T × d} (d = model dim), project to Q, K, V:
Q = X W_Q, K = X W_K, V = X W_V ∈ ℝ^{T × n_heads × head_dim}Attention output:
Y = softmax(Q Kᵀ / √head_dim) V ∈ ℝ^{T × n_heads × head_dim}Complexity per layer: O(T² · head_dim · n_heads) in time, O(T · n_heads · head_dim) in KV cache size per token.
MQA — n_kv_groups = 1#
K and V are projected with a single shared head; Q keeps n_heads. The KV cache size per token shrinks by n_heads, but the attention math is unchanged.
GQA — n_kv_groups ∈ (1, n_heads)#
K and V are projected to n_kv_groups heads; Q is projected to n_heads heads; each query head attends to one KV group (round-robin assignment: query head i reads from KV group i // (n_heads // n_kv_groups)).
KV cache size per token: n_heads / n_kv_groups × MHA.
MLA — DeepSeek-style#
Project X once into a shared low-rank latent c_t:
c_t = X · W_KV_A ∈ ℝ^{T × kv_lora_rank}
k_t = c_t · W_K_B ∈ ℝ^{T × n_heads × qk_nope_head_dim}
v_t = c_t · W_V_B ∈ ℝ^{T × n_heads × v_head_dim}Each query also has its own low-rank projection:
q_nope_t = X · W_Q_A · W_Q_B ∈ ℝ^{T × n_heads × qk_nope_head_dim}The "absorbed" trick: precompute W_K_B and W_V_B once, then absorb them into the query projection, so the attention op sees a single shared latent c_t plus per-head q_t. The KV cache becomes c_t only — T × kv_lora_rank — not T × n_heads × (qk_nope_head_dim + v_head_dim).
This is the innovation of MLA: the cache shrinks by another factor of n_heads × (qk_nope + v_head_dim) / kv_lora_rank (roughly n_heads × 2 × head_dim / kv_lora_rank).
HyMo's choice — MQA-4#
| Knob | Value |
|---|---|
n_heads | 16 |
n_kv_groups | 4 |
q_lora_rank | 224 |
kv_lora_rank | 128 |
head_dim | 128 |
qk_rope_head_dim | 32 |
qk_nope_head_dim | 96 |
v_head_dim | 128 |
These are not full MLA — there's no W_K_B/W_V_B absorption. Instead, every query head reads from one of 4 KV heads (round-robin: query heads {0,1,2,3} → KV head 0; query heads {4,5,6,7} → KV head 1; …).
The q_lora_rank = 224 and kv_lora_rank = 128 are full MLA: the query is first projected to 224 dim, then up-projected to n_heads × qk_nope_head_dim; the KV is projected to a shared 128-dim latent per token and broadcast to the 4 KV heads. The KV-cache win is from kv_lora_rank = 128 instead of 4 × 128 = 512.
So HyMo is "MQA-4 with low-rank KV compression" — MQA-style KV sharing plus the MLA-style latent. Pure MLA absorption (the trick that lets you store only c_t) is what was deferred; at this scale, the simpler kernel wins.
Implementation in HyMo#
src/hymo/models/mla.py:MultiHeadLatentAttention—class MultiHeadLatentAttention
(the projection class).src/hymo/models/mla.py:MultiHeadLatentAttention.__init__—__init__with the 8 MLA config
fields used.src/hymo/models/mla.py:MultiHeadLatentAttention.n_heads— propertiesn_heads,
n_kv_groups,qk_rope_head_dim,qk_nope_head_dim,v_head_dim.src/hymo/models/mla.py:MultiHeadLatentAttention.forward—forward(x): low-rank Q → split
into RoPE / NoPE parts; KV → 4 groups with partial RoPE on the firstqk_rope_head_dim = 32ofhead_dim; attention; output projection.src/hymo/models/mla.py:MLABlock—class MLABlock: pre-norm +
MultiHeadLatentAttention + residual.src/hymo/models/mla.py:MLABlock.__init__—__init__(config, layer_idx):
builds the block (pre-norm + attention + MoE + residuals).src/hymo/models/mla.py:MLABlock.forward—forward(x): full block
forward, including the soft-cap (no — the softcap is on the logits, not the attention output; seemodel.py:HyMo.softcap).
Worked example#
HyMo at production scale:
- 16 query heads × 128 head_dim = 2048 query projection width.
- 4 KV heads × 128 = 512 KV projection width (or via latent:
128-dim shared latent broadcast to 4 heads). - Attention FLOPs per layer per token:
2 * (q_dim * k_dim + q_dim * v_dim) ≈ 2 * (2048 * 512) = 2.1 MFLOPs, plus theO(T²)softmax matmul. - 8 MLA layers per forward,
T = 4096, batch 4: ~`2.1 M 4096
8 = 69 G` FLOPs for QK + softmax.
The shared latent saves ~4× on KV cache size: from (4096 * 16 * 256 * 2) (MHA, both K and V) bytes down to (4096 * 128 * 2) = 1 MB per token — small but compounded across the 8 MLA layers at inference time.
Interview Q&A#
Q1. Why does MLA beat MQA at long context?
A: MQA collapses all KV heads into one; MLA collapses all KV heads into one plus projects through a low-rank bottleneck that drops the head dim entirely. For very long contexts (T > 32 k), the bottleneck gives another factor of ~head_dim / kv_lora_rankin cache size, which matters for inference memory.
Q2. Why didn't HyMo go full MLA?
A: Two reasons. First, at 434M active params the absorbed-MLA kernel (which needs extra matmuls to project through the low-rank bottleneck at attention time) is not a clear win on quality-per-FLOPs vs. MQA-4. Second, the absorbed-MLA kernel is harder to write in pure PyTorch + Triton than the MQA-4 kernel; HyMo's kernel budget didn't include it. The kv_lora_rank = 128 field is set up so a future v1.2 could flip to full MLA absorption without changing the config shape.
Q3. Why is the partial RoPE only on the first 25% of head_dim?
A: Empirically, RoPE on the fullhead_dimworks but is over-parameterized; 25% gives the position-aware channels enough rope to learn relative-position patterns without burning 75% of the head dim on position encoding (which would otherwise compete with content attention). Seemodel-architecture.md.
Q4. What does "absorbed" mean in MLA?
A: The trick that the latentc_tcan be the only thing stored in the KV cache (instead of per-headk_t,v_t). The per-head projectionsW_K_B,W_V_Bare folded into the query projection at training time, so at inference, attention readsc_tdirectly and the per-headk_t,v_tare reconstructed on the fly. HyMo doesn't do this folding yet; seeoptimization.mdfor the future work.
Q5. MQA-4 vs GQA-1.75 — what's the difference?
A: GQA-1.75 (an earlier draft of HyMo's config) means the ration_heads / n_kv_groups = 1.75— but1.75isn't an integer, so it was a placeholder for "not yet decided". MQA-4 isn_kv_groups = 4, so the ratio is16 / 4 = 4. Ablation family D (D_mqa4_vs_gqa175) compares these.
Cross-links#
- [
model-architecture.md(model-architecture.md) §4 — the
MLA block walkthrough. concepts/model-architecture.md—
partial RoPE on the first 25% ofhead_dim.concepts/gdn-and-mla.md—
why MLA + GDN, not just MLA.
Learning objectives#
After this file, you can:
- State Rotary Position Embeddings (RoPE) and how they preserve
relative position through a rotation. - Explain HyMo's "partial RoPE" choice (25% of
head_dim). - Describe the NoPE-hybrid ablation (no PE on select GDN
layers) and why it's deferred to v1.1.
Intuition#
The attention score between query q_i and key k_j depends on the content of q_i and k_j. Without position information, the model cannot distinguish "the cat sat on the mat" from "mat the on sat cat the" — same tokens, different order.
Two main approaches:
| Approach | What it stores | Compute cost |
|---|---|---|
| Absolute position embeddings | A learned vector per position, added to the token embedding. | One extra vocab_size + max_seq_len embedding table. |
| RoPE (Su et al. 2021) | A learned rotation per head, applied to q and k per position. | Per-position rotation; no extra parameters. |
| NoPE (no position encoding) | None — the model has no explicit position signal. | Zero overhead. |
RoPE is the modern default. It encodes relative position through an angle: at position p, query/key vectors are rotated by an angle p · θ_k where θ_k is a per-dim frequency. The attention score after rotation depends only on the relative position (i - j), not on the absolute positions.
Partial RoPE: instead of applying RoPE to the full head_dim, apply it to only the first qk_rope_head_dim of it. The remaining qk_nope_head_dim = head_dim - qk_rope_head_dim carries no position information.
# HyMo at production scale:
# qk_rope_head_dim = 32
# qk_nope_head_dim = 96
# head_dim = 128 (= 32 + 96)
# 25% RoPE, 75% NoPE per headWhy 25%?#
Empirically (DeepSeek-V2, V3; also some Llama ablation work): applying RoPE to fewer dimensions improves long-context performance. The intuition: the head-dim budget is shared between content and position; giving 75% of it to content lets each head specialize more. The 25% is enough rope to learn relative position patterns.
The NoPE-hybrid (deferred)#
An ablation that disables position encoding on select GDN layers — specifically, the GDN layer immediately after each MLA layer. The 7 affected positions for an 8-MLA stack are {3, 7, 11, 15, 19, 23, 27}.
Why? A GDN layer with no position information relies entirely on its recurrence to track "where" tokens are. The MLA-then-GDN sandwich is a natural place to test this — the MLA layer has its own position info (rotated by partial RoPE), and the GDN that follows can either reuse it (via the hidden state) or have no position info at all. The ablation family D_mqa4_vs_gqa175 and the NoPE-hybrid flag (nope_hybrid_gdn_enabled: false in v1.0) are what gate this.
The risk was that "no position" hurts the model. The mitigation in v1.0 is simple: ship with the flag off (everyone gets partial RoPE), defer the test to v1.1.
Math derivation#
RoPE#
A 2D rotation by angle θ:
R(θ) = [cos θ, -sin θ]
[sin θ, cos θ]For a head_dim-dim vector, RoPE applies a different angle to each pair of dimensions. Standard convention:
θ_k = 1 / (rope_theta^(2k / head_dim)) k = 0, 1, ..., head_dim/2 - 1At position p, the rotation per pair is p · θ_k. Query q at position p:
q_rot = RoPE(q, p) = [q_0 cos(p θ_0) - q_1 sin(p θ_0),
q_0 sin(p θ_0) + q_1 cos(p θ_0),
q_2 cos(p θ_1) - q_3 sin(p θ_1),
...]Same for k at position j. The attention inner product q_rot · k_rot depends on (p - j) (modulo 2π), so the model learns relative position patterns.
Partial RoPE#
Apply RoPE only to the first qk_rope_head_dim of q and k:
q_rot = [RoPE(q[:32], p), q[32:]] # 32 dims rotated, 96 dims not
k_rot = [RoPE(k[:32], j), k[32:]]The remaining qk_nope_head_dim = 96 dims are vanilla content attention. The attention inner product is:
q_rot · k_rot = RoPE(q[:32], p) · RoPE(k[:32], j)
+ q[32:128] · k[32:128]The first term has relative-position dependence; the second doesn't. This mixing is what gives partial RoPE its quality advantage over full RoPE.
NoPE#
Set qk_rope_head_dim = 0. The full head_dim carries no position information. The model must rely on:
- The previous MLA layer's hidden state (which has its own
partial RoPE). - The recurrence state of the GDN (which carries implicit
position through the cumulative decay). - The MoE router (which sees the same hidden state).
Empirically, partial NoPE-hybrid (mixing RoPE and NoPE across layers) is competitive with full RoPE on most tasks but better on long-context benchmarks. The v1.1 ablation will test this.
Implementation in HyMo#
src/hymo/models/rope.py:RotaryEmbedding—class RotaryEmbedding.src/hymo/models/rope.py:RotaryEmbedding.__init__—__init__: precomputes the
cos,sintables formax_seq_lenpositions andhead_dim / 2freq pairs.src/hymo/models/rope.py:RotaryEmbedding.apply_rope— `apply_rope(x, positions, *,
start_pos=0): applies the rotation to the firstqk_rope_head_dimofx`; leaves the rest unchanged.src/hymo/models/rope.py:RotaryEmbedding.extra_repr—extra_repr.
Wiring:
src/hymo/models/mla.py:MultiHeadLatentAttention.forward—forward(x)calls
apply_rope(q_rope, positions)for the firstqk_rope_head_dim = 32ofqandk.src/hymo/models/gdn.py:GatedDeltaNetBlock.forward—forward(x)calls
apply_ropeon the GDN'sbandckeys (which have the same shape as MLA'sq_rope).src/hymo/models/model.py:HyMo.__init__— for each layer, if
i in nope_hybrid_gdn_positions, the GDN block is built withuse_rope=False(noapply_ropecall).
The use_rope flag on GatedDeltaNetBlock is set in model.py:HyMo.__init__:
for i in range(config.n_layers):
if i in mla_positions:
self.layers.append(MLABlock(config, layer_idx=i))
else:
use_rope = i not in nope_hybrid
self.layers.append(
GatedDeltaNetBlock(config, layer_idx=i, use_rope=use_rope)
)So nope_hybrid_gdn_enabled: true means the 7 GDN layers at positions {3, 7, 11, 15, 19, 23, 27} are GDN-without-RoPE, the other 17 GDN layers get partial RoPE.
Worked example#
Production scale (configs/hymo_750m.yaml):
head_dim = 128,qk_rope_head_dim = 32,qk_nope_head_dim = 96rope_theta = 10_000.0- 32 layers, 8 MLA + 24 GDN
nope_hybrid_gdn_enabled: false(v1.0 default)
Per-layer positional info:
| Layer | Type | Pos info |
|---|---|---|
| 0 | MLA | 32-dim RoPE on q, k |
| 1, 2 | GDN | 32-dim RoPE on b, c |
| 3 | GDN | 32-dim RoPE on b, c (NoPE would be here if nope_hybrid_gdn_enabled: true) |
| 4 | MLA | 32-dim RoPE on q, k |
| ... | ... | ... |
| 28 | MLA | 32-dim RoPE on q, k |
| 29, 30, 31 | GDN | 32-dim RoPE on b, c |
If the NoPE-hybrid flag were flipped, layers 3, 7, 11, 15, 19, 23, 27 would not apply RoPE — the rest of the model looks identical.
Interview Q&A#
Q1. Why partial RoPE (25%) instead of full RoPE?
A: Empirically, full RoPE on 128-dim heads over-allocates capacity to position encoding. 25% gives enough rotational channels to learn relative positions while leaving 75% of the head dim for content signal. This is the DeepSeek-V2/V3 choice.
Q2. Why does the NoPE-hybrid defer to v1.1?
A: Risk reduction. The shipped model is the primary 30 B pre-training run; flipping the NoPE-hybrid flag without a prior ablation result is a multi-day experiment with no guarantee of payoff. The v1.0 ships with the flag off (everyone gets partial RoPE); the v1.1 ablation tests whether the 7 NoPE GDN layers help long-context tasks.
Q3. Why is rope_theta = 10_000 and not 500_000 (Llama-3)?
A: 10_000 is the original RoPE default; it gives reasonable extrapolation up to ~32 k context. 500_000 is the Llama-3 "extended context" choice, which extrapolates to 128 k but adds no quality at the trained context. HyMo trains at max_seq_len = 4_096; 10_000 is sufficient for that range without wasted parameter capacity.
Q4. Why apply RoPE to b and c in GDN, not just to q and k in MLA?
A: The GDN writes-to-state viaband reads-from-state viac. Ifbandcdon't carry position information, the model can't write a "memory at position p" or read "memory from position j" — the recurrence becomes position-agnostic, which hurts associative recall. Applying RoPE tob, clets the GDN track timing through the recurrence.
Q5. Why does apply_rope take a start_pos argument?
A: For inference with KV-cache reuse. When decoding at positionp, the model only needs the rotation atp, not all positions.start_poslets the kernel slice the cos/sin tables without recomputing. Currentlystart_pos=0always (no incremental inference in v1.0), but the argument is there for the v1.1 inference refactor.
Q6. What's the cost of partial RoPE vs. full RoPE?
A:apply_ropealways slices and rotates the firstqk_rope_head_dimof the input. Whether that's 32 or 128 dims, the cost is one rotation per dim pair (~30 ops per pair). The 75% NoPE portion is "free" — no rotation compute. Net: partial RoPE is cheaper than full RoPE.
Cross-links#
model-architecture.md§7
(RoPE walkthrough).model-architecture.md— how RoPE is
applied to MLA'sqandk.gdn-and-mla.md—
how RoPE is applied to GDN'sbandc.gdn-and-mla.md—
the NoPE-hybrid as an ablation.
References#
- gdn-and-mla.md — the GDN/MLA/MoE/MTP mechanism deep-dives.
- optimization.md — optimizers, scheduler, FSDP-2, init status.
- kernels.md — the Triton GDN kernel.
- design.md — the full architecture & design document.
- training.md — the training pipeline.
- config.md — the config system reference.
- Source:
src/hymo/models/model.py,src/hymo/models/mla.py,src/hymo/models/gdn.py,src/hymo/models/rope.py,src/hymo/models/moe.py,src/hymo/models/mtp.py,src/hymo/models/gdn_triton.py.
Next: Read training.md for the data loading, tokenization, and training pipeline.