mid-run checkpoint step 8000 (0.26B tokens)
Browse files- README.md +100 -0
- model.safetensors +3 -0
- pssa_kernels.py +151 -0
- pssa_model.py +179 -0
- tokenizer.json +0 -0
README.md
ADDED
|
@@ -0,0 +1,100 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
---
|
| 2 |
+
license: gpl-3.0
|
| 3 |
+
library_name: model
|
| 4 |
+
tags:
|
| 5 |
+
- state-space-model
|
| 6 |
+
- recurrent
|
| 7 |
+
- language-model
|
| 8 |
+
- from-scratch
|
| 9 |
+
- plastic
|
| 10 |
+
- mamba-like
|
| 11 |
+
- triton
|
| 12 |
+
datasets:
|
| 13 |
+
- roneneldan/TinyStories
|
| 14 |
+
- HuggingFaceFW/fineweb-edu
|
| 15 |
+
metrics:
|
| 16 |
+
- perplexity
|
| 17 |
+
- accuracy
|
| 18 |
+
---
|
| 19 |
+
|
| 20 |
+
# pssa-100m
|
| 21 |
+
|
| 22 |
+
A 109M-parameter language model with no attention, trained from scratch in a live notebook on one RTX PRO 6000 Blackwell. The [plastic state-space architecture](https://github.com/Sparticle62ops/pssa) (PSSA: selective diagonal state-space scan, Poincaré-ball episodic memory read, gated adapter, SiLU MLP) reimplemented in PyTorch, so the architecture's first 100M-scale checkpoint exists as plain weights you can load and sample in a few lines.
|
| 23 |
+
|
| 24 |
+
The Rust repo's headline result is a 1.5M-parameter PSSA beating a matched transformer by 0.45 nats on WikiText-103. This checkpoint is the architecture's first answer at 109M: 2.5B tokens in three runs on one GPU, loss curves and sample text below. What you get is a toy that writes real English in two registers. What you don't get is facts. At 109M params it has capacity for style, not knowledge, and the web text it emits is confidently invented.
|
| 25 |
+
|
| 26 |
+
## Architecture
|
| 27 |
+
|
| 28 |
+
18 stacked blocks of:
|
| 29 |
+
|
| 30 |
+
```
|
| 31 |
+
xn = rmsnorm(x)
|
| 32 |
+
delta = softplus(W_delta xn) per-channel step size
|
| 33 |
+
h_t = exp(delta_t·A) h_{t-1} + delta_t·B_t·x_t diagonal SSM scan, A < 0
|
| 34 |
+
y_t = <h_t, C_t>
|
| 35 |
+
q = poincare(W_qx xn + W_qh y) query from token AND state
|
| 36 |
+
w = softmax(-d_H(q, k)/tau) hyperbolic distance to 512 memory slots
|
| 37 |
+
m = w @ v bounded episodic read
|
| 38 |
+
z = y/√d_s + gate·W_proj m + rank16-silu-adapter(xn) + MLP
|
| 39 |
+
h_i = h_{i-1} + block_i(h)/√depth
|
| 40 |
+
```
|
| 41 |
+
|
| 42 |
+
Parameter counts: 101.50M core (d_model 768, untied 16k-token embedding and head) plus 7.37M across the 18 memory banks. There is no attention anywhere in the model. Cost per token is constant in context length; the state is fixed-size.
|
| 43 |
+
|
| 44 |
+
Two deviations from the Rust reference, both forced by scale: the memory bank is gradient-trained rather than surprise-gated (the reference's 3.5-nat novelty threshold stops firing within ~300 steps once the model is past the basics, so the write path would sit inert for the rest of a 2.5B-token run), and backprop is truncated at 256 tokens. Training used a fused Triton scan (forward + backward in one kernel each) that matches a plain PyTorch loop to ~1e-7 on every gradient.
|
| 45 |
+
|
| 46 |
+
## Training
|
| 47 |
+
|
| 48 |
+
Three runs, one GPU, always resuming the same checkpoint:
|
| 49 |
+
|
| 50 |
+
| run | data | tokens | speed |
|
| 51 |
+
|---|---|---|---|
|
| 52 |
+
| 1 | TinyStories only | 400M | 71k tok/s |
|
| 53 |
+
| 2 | 50/50 TinyStories / FineWeb-Edu | +1.1B | 128k tok/s |
|
| 54 |
+
| 3 | fine-tune, 85% FineWeb-Edu | +1.0B | 128k tok/s |
|
| 55 |
+
|
| 56 |
+
The speed jump between runs 1 and 2 is `torch.compile` on the block: Inductor fuses ~40 small elementwise ops (the Poincaré distances, memory softmax, gates) into a handful of kernels, 184→130 ms/step, gradient cosine 0.9998 vs eager. Batch went 32→64 and chunk 128→256 on top of that. 99% GPU utilization held for the whole of runs 2 and 3.
|
| 57 |
+
|
| 58 |
+
## Results
|
| 59 |
+
|
| 60 |
+
Held-out, 1M unseen tokens per corpus:
|
| 61 |
+
|
| 62 |
+
| | cross-entropy | perplexity | next-token acc |
|
| 63 |
+
|---|---|---|---|
|
| 64 |
+
| TinyStories | 1.329 | 3.78 | 64.5% |
|
| 65 |
+
| FineWeb-Edu web | 2.935 | 18.8 | 43.4% |
|
| 66 |
+
|
| 67 |
+
The interesting datapoint is run 2: moving from story-only to a 50/50 mixture *improved* story loss (1.302 → 1.244) while learning web text at all. Story loss only gave back ground when the mixture went to 85% web.
|
| 68 |
+
|
| 69 |
+
Samples, temp 0.7 top-k 40:
|
| 70 |
+
|
| 71 |
+
> Once upon a time there was a tiny dragon. The dragon was very fast and loved to run. He had a big, soft voice and a long tail. He lived in a deep, dark cave. One day, the dragon heard a voice. It was a little bird. The bird was lost and scared...
|
| 72 |
+
|
| 73 |
+
> In the summer of 1969, the U.S. Department of Agriculture announced that it would implement "universally mandated standards for the development of new products and products." In the first phase of its series of "Design, Material and Technologies", the Standards for Energy and Smart Buildings (SEP) was developed and funded...
|
| 74 |
+
|
| 75 |
+
Web loss was still falling when the budget ran out. This model is underfit, not overfit; more tokens would keep helping.
|
| 76 |
+
|
| 77 |
+
## Use
|
| 78 |
+
|
| 79 |
+
```python
|
| 80 |
+
# pip install torch safetensors tokenizers
|
| 81 |
+
from pssa_model import from_pretrained, generate
|
| 82 |
+
|
| 83 |
+
model, tok = from_pretrained(".") # dir containing model.safetensors + tokenizer.json
|
| 84 |
+
print(generate(model, tok, "Once upon a time", max_new=100))
|
| 85 |
+
```
|
| 86 |
+
|
| 87 |
+
`pssa_model.py` is self-contained: `PSSAConfig`, `PSSAModel`, loader, sampler. On CPU it runs a plain PyTorch scan; on CUDA it uses the fused Triton kernel from `pssa_kernels.py` when that file is importable. No transformers, no configuration — plain weights and ~150 lines.
|
| 88 |
+
|
| 89 |
+
Weights are bf16 safetensors, 217 MB. The 16,384-token byte-level BPE tokenizer trained on TinyStories ships alongside. It handles general web text without unmangling, at some cost in efficiency on rare Unicode.
|
| 90 |
+
|
| 91 |
+
## Files
|
| 92 |
+
|
| 93 |
+
- `model.safetensors`: weights, bf16
|
| 94 |
+
- `pssa_model.py`: model + loader + sampler
|
| 95 |
+
- `pssa_kernels.py`: fused Triton scan
|
| 96 |
+
- `tokenizer.json`: 16k byte-level BPE
|
| 97 |
+
|
| 98 |
+
## Credit
|
| 99 |
+
|
| 100 |
+
Architecture, reference implementation and the small-scale comparison: [Sparticle62ops/pssa](https://github.com/Sparticle62ops/pssa) (GPL-3.0). This checkpoint was trained in a marimo notebook on a single RTX PRO 6000 Blackwell; the training notebook is in the repo's `notebooks/` directory.
|
model.safetensors
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:751607c8f8f2dc9d30a431b902ca20d5219407d429046811af493d451ada567b
|
| 3 |
+
size 217776144
|
pssa_kernels.py
ADDED
|
@@ -0,0 +1,151 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
"""Fused Triton scan for the PSSA selective diagonal state-space recurrence.
|
| 2 |
+
|
| 3 |
+
h_t = exp(delta_t * A) * h_{t-1} + delta_t * b_t * x_t
|
| 4 |
+
y_t = <h_t, c_t>
|
| 5 |
+
|
| 6 |
+
`A` is (D, DS) negative, `delta` is (B, T, D) positive, `b`/`c` are (B, T, DS).
|
| 7 |
+
The whole (forward + backward) runs as one kernel launch per chunk, with the
|
| 8 |
+
per-token state `h_t` materialised so the backward pass is a single reverse
|
| 9 |
+
sweep. Carries `h0` across chunks and returns `hT`.
|
| 10 |
+
"""
|
| 11 |
+
import torch
|
| 12 |
+
import triton
|
| 13 |
+
import triton.language as tl
|
| 14 |
+
|
| 15 |
+
__all__ = ["ScanFn", "scan_reference", "selftest"]
|
| 16 |
+
|
| 17 |
+
|
| 18 |
+
@triton.jit
|
| 19 |
+
def scan_fwd(x_ptr, d_ptr, bp_ptr, cp_ptr, A_ptr, h0_ptr, y_ptr, hT_ptr, H_ptr,
|
| 20 |
+
T: tl.constexpr, D: tl.constexpr, DS: tl.constexpr, BD: tl.constexpr):
|
| 21 |
+
b = tl.program_id(0)
|
| 22 |
+
pd = tl.program_id(1)
|
| 23 |
+
od = pd * BD + tl.arange(0, BD)
|
| 24 |
+
os_ = tl.arange(0, DS)
|
| 25 |
+
A = tl.load(A_ptr + od[:, None] * DS + os_[None, :])
|
| 26 |
+
h = tl.load(h0_ptr + b * D * DS + od[:, None] * DS + os_[None, :])
|
| 27 |
+
for t in range(T):
|
| 28 |
+
x = tl.load(x_ptr + b * T * D + t * D + od).to(tl.float32)
|
| 29 |
+
dl = tl.load(d_ptr + b * T * D + t * D + od).to(tl.float32)
|
| 30 |
+
bp = tl.load(bp_ptr + b * T * DS + t * DS + os_).to(tl.float32)
|
| 31 |
+
cp = tl.load(cp_ptr + b * T * DS + t * DS + os_).to(tl.float32)
|
| 32 |
+
h = tl.exp(dl[:, None] * A) * h + dl[:, None] * bp[None, :] * x[:, None]
|
| 33 |
+
tl.store(y_ptr + b * T * D + t * D + od, tl.sum(h * cp[None, :], axis=1))
|
| 34 |
+
tl.store(H_ptr + b * T * D * DS + t * D * DS + od[:, None] * DS + os_[None, :], h)
|
| 35 |
+
tl.store(hT_ptr + b * D * DS + od[:, None] * DS + os_[None, :], h)
|
| 36 |
+
|
| 37 |
+
|
| 38 |
+
@triton.jit
|
| 39 |
+
def scan_bwd(dy_ptr, x_ptr, d_ptr, bp_ptr, cp_ptr, A_ptr, H_ptr, h0_ptr,
|
| 40 |
+
gh0_ptr, gx_ptr, gd_ptr, gbp_ptr, gcp_ptr, gA_ptr,
|
| 41 |
+
T: tl.constexpr, D: tl.constexpr, DS: tl.constexpr, BD: tl.constexpr):
|
| 42 |
+
b = tl.program_id(0)
|
| 43 |
+
pd = tl.program_id(1)
|
| 44 |
+
od = pd * BD + tl.arange(0, BD)
|
| 45 |
+
os_ = tl.arange(0, DS)
|
| 46 |
+
A = tl.load(A_ptr + od[:, None] * DS + os_[None, :])
|
| 47 |
+
g = tl.zeros((BD, DS), tl.float32)
|
| 48 |
+
gA = tl.zeros((BD, DS), tl.float32)
|
| 49 |
+
for t in range(T - 1, -1, -1):
|
| 50 |
+
dy = tl.load(dy_ptr + b * T * D + t * D + od).to(tl.float32)
|
| 51 |
+
x = tl.load(x_ptr + b * T * D + t * D + od).to(tl.float32)
|
| 52 |
+
dl = tl.load(d_ptr + b * T * D + t * D + od).to(tl.float32)
|
| 53 |
+
bp = tl.load(bp_ptr + b * T * DS + t * DS + os_).to(tl.float32)
|
| 54 |
+
cp = tl.load(cp_ptr + b * T * DS + t * DS + os_).to(tl.float32)
|
| 55 |
+
h = tl.load(H_ptr + b * T * D * DS + t * D * DS + od[:, None] * DS + os_[None, :])
|
| 56 |
+
if t < T - 1:
|
| 57 |
+
dln = tl.load(d_ptr + b * T * D + (t + 1) * D + od).to(tl.float32)
|
| 58 |
+
g = dy[:, None] * cp[None, :] + tl.exp(dln[:, None] * A) * g
|
| 59 |
+
else:
|
| 60 |
+
g = dy[:, None] * cp[None, :]
|
| 61 |
+
if t > 0:
|
| 62 |
+
hp = tl.load(H_ptr + b * T * D * DS + (t - 1) * D * DS + od[:, None] * DS + os_[None, :])
|
| 63 |
+
else:
|
| 64 |
+
hp = tl.load(h0_ptr + b * D * DS + od[:, None] * DS + os_[None, :])
|
| 65 |
+
a = tl.exp(dl[:, None] * A)
|
| 66 |
+
ga = g * hp
|
| 67 |
+
gd = tl.sum(ga * a * A, axis=1) + tl.sum(g * bp[None, :] * x[:, None], axis=1)
|
| 68 |
+
gbp = tl.sum(g * dl[:, None] * x[:, None], axis=0)
|
| 69 |
+
gcp = tl.sum(dy[:, None] * h, axis=0)
|
| 70 |
+
gx = tl.sum(g * dl[:, None] * bp[None, :], axis=1)
|
| 71 |
+
gA += ga * a * dl[:, None]
|
| 72 |
+
if t == 0:
|
| 73 |
+
tl.store(gh0_ptr + b * D * DS + od[:, None] * DS + os_[None, :], a * g)
|
| 74 |
+
tl.store(gx_ptr + b * T * D + t * D + od, gx)
|
| 75 |
+
tl.store(gd_ptr + b * T * D + t * D + od, gd)
|
| 76 |
+
tl.atomic_add(gbp_ptr + b * T * DS + t * DS + os_, gbp)
|
| 77 |
+
tl.atomic_add(gcp_ptr + b * T * DS + t * DS + os_, gcp)
|
| 78 |
+
tl.atomic_add(gA_ptr + od[:, None] * DS + os_[None, :], gA)
|
| 79 |
+
|
| 80 |
+
|
| 81 |
+
class ScanFn(torch.autograd.Function):
|
| 82 |
+
"""Differentiable fused scan. Returns (y, hT); hT is not differentiated."""
|
| 83 |
+
|
| 84 |
+
@staticmethod
|
| 85 |
+
def forward(ctx, x, delta, bp, cp, A, h0, bd=16):
|
| 86 |
+
B, T, D = x.shape
|
| 87 |
+
DS = A.shape[1]
|
| 88 |
+
y = torch.empty(B, T, D, device=x.device, dtype=torch.float32)
|
| 89 |
+
hT = torch.empty(B, D, DS, device=x.device, dtype=torch.float32)
|
| 90 |
+
H = torch.empty(B, T, D, DS, device=x.device, dtype=torch.float32)
|
| 91 |
+
scan_fwd[(B, triton.cdiv(D, bd))](
|
| 92 |
+
x, delta, bp, cp, A, h0, y, hT, H, T=T, D=D, DS=DS, BD=bd)
|
| 93 |
+
ctx.save_for_backward(x, delta, bp, cp, A, H, h0)
|
| 94 |
+
ctx.bd = bd
|
| 95 |
+
ctx.mark_non_differentiable(hT)
|
| 96 |
+
return y, hT
|
| 97 |
+
|
| 98 |
+
@staticmethod
|
| 99 |
+
def backward(ctx, dy, dhT):
|
| 100 |
+
x, delta, bp, cp, A, H, h0 = ctx.saved_tensors
|
| 101 |
+
B, T, D = x.shape
|
| 102 |
+
DS = A.shape[1]
|
| 103 |
+
bd = ctx.bd
|
| 104 |
+
gx = torch.empty(B, T, D, device=x.device, dtype=torch.float32)
|
| 105 |
+
gd = torch.empty_like(delta, dtype=torch.float32)
|
| 106 |
+
gbp = torch.zeros_like(bp, dtype=torch.float32)
|
| 107 |
+
gcp = torch.zeros_like(cp, dtype=torch.float32)
|
| 108 |
+
gA = torch.zeros_like(A, dtype=torch.float32)
|
| 109 |
+
gh0 = torch.zeros(B, D, DS, device=x.device, dtype=torch.float32)
|
| 110 |
+
scan_bwd[(B, triton.cdiv(D, bd))](
|
| 111 |
+
dy.contiguous(), x, delta, bp, cp, A, H, h0, gh0,
|
| 112 |
+
gx, gd, gbp, gcp, gA, T=T, D=D, DS=DS, BD=bd)
|
| 113 |
+
return gx.to(x.dtype), gd, gbp, gcp, gA, gh0.to(h0.dtype), None
|
| 114 |
+
|
| 115 |
+
|
| 116 |
+
def scan_reference(x, delta, bp, cp, A, h0):
|
| 117 |
+
"""Plain PyTorch loop, for tests. Returns (y, H, hT)."""
|
| 118 |
+
B, T, D = x.shape
|
| 119 |
+
h = h0
|
| 120 |
+
ys, hs = [], []
|
| 121 |
+
for t in range(T):
|
| 122 |
+
h = torch.exp(delta[:, t, :, None] * A) * h + delta[:, t, :, None] * bp[:, t, None, :] * x[:, t, :, None]
|
| 123 |
+
ys.append((h * cp[:, t, None, :]).sum(-1))
|
| 124 |
+
hs.append(h)
|
| 125 |
+
return torch.stack(ys, 1), torch.stack(hs, 1), h
|
| 126 |
+
|
| 127 |
+
|
| 128 |
+
def selftest(ds=16, d=64, b=2, t=7, verbose=True):
|
| 129 |
+
torch.manual_seed(1)
|
| 130 |
+
dev = "cuda"
|
| 131 |
+
x = torch.randn(b, t, d, device=dev, requires_grad=True)
|
| 132 |
+
delta = torch.nn.functional.softplus(torch.randn(b, t, d, device=dev)).requires_grad_(True)
|
| 133 |
+
bp = torch.randn(b, t, ds, device=dev, requires_grad=True)
|
| 134 |
+
cp = torch.randn(b, t, ds, device=dev, requires_grad=True)
|
| 135 |
+
A = (-torch.nn.functional.softplus(torch.randn(d, ds, device=dev))).requires_grad_(True)
|
| 136 |
+
h0 = (torch.randn(b, d, ds, device=dev) * 0.5).requires_grad_(True)
|
| 137 |
+
yr, _, _ = scan_reference(x, delta, bp, cp, A, h0)
|
| 138 |
+
y, hT = ScanFn.apply(x, delta, bp, cp, A, h0)
|
| 139 |
+
gy = torch.randn_like(yr)
|
| 140 |
+
gref = torch.autograd.grad(yr, (x, delta, bp, cp, A, h0), gy)
|
| 141 |
+
gtri = torch.autograd.grad(y, (x, delta, bp, cp, A, h0), gy)
|
| 142 |
+
out = {"fwd": ((y - yr).abs().max() / yr.abs().max()).item()}
|
| 143 |
+
for nm, a, bb in zip(["dx", "dd", "dbp", "dcp", "dA", "dh0"], gtri, gref):
|
| 144 |
+
out[nm] = ((a - bb).abs().max() / bb.abs().max()).item()
|
| 145 |
+
if verbose:
|
| 146 |
+
print("scan selftest rel errs: " + " ".join(f"{k}={v:.1e}" for k, v in out.items()))
|
| 147 |
+
return out
|
| 148 |
+
|
| 149 |
+
|
| 150 |
+
if __name__ == "__main__":
|
| 151 |
+
selftest()
|
pssa_model.py
ADDED
|
@@ -0,0 +1,179 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
"""Standalone PSSA model for the Hub: numpy-only inference, bf16 safetensors."""
|
| 2 |
+
import math
|
| 3 |
+
from dataclasses import dataclass
|
| 4 |
+
|
| 5 |
+
import numpy as np
|
| 6 |
+
import torch
|
| 7 |
+
from torch import nn
|
| 8 |
+
import torch.nn.functional as F
|
| 9 |
+
|
| 10 |
+
try:
|
| 11 |
+
from pssa_kernels import ScanFn
|
| 12 |
+
except Exception: # no compiled kernel (CPU-only install): fall back below
|
| 13 |
+
ScanFn = None
|
| 14 |
+
|
| 15 |
+
|
| 16 |
+
def _scan(x, delta, bp, cp, A, h0):
|
| 17 |
+
"""h_t = exp(delta_t·A) h_{t-1} + delta_t·b_t·x_t, y_t = <h_t, c_t>."""
|
| 18 |
+
if ScanFn is not None and x.is_cuda:
|
| 19 |
+
return ScanFn.apply(x, delta, bp, cp, A, h0)
|
| 20 |
+
B, T, D = x.shape
|
| 21 |
+
h = h0
|
| 22 |
+
ys = []
|
| 23 |
+
for t in range(T):
|
| 24 |
+
h = torch.exp(delta[:, t, :, None] * A) * h + delta[:, t, :, None] * bp[:, t, None, :] * x[:, t, :, None]
|
| 25 |
+
ys.append((h * cp[:, t, None, :]).sum(-1))
|
| 26 |
+
return torch.stack(ys, 1), h
|
| 27 |
+
|
| 28 |
+
|
| 29 |
+
@dataclass
|
| 30 |
+
class PSSAConfig:
|
| 31 |
+
vocab_size: int = 16384
|
| 32 |
+
d_model: int = 768
|
| 33 |
+
depth: int = 18
|
| 34 |
+
d_state: int = 16
|
| 35 |
+
d_key: int = 32
|
| 36 |
+
mem_slots: int = 512
|
| 37 |
+
adapter_rank: int = 16
|
| 38 |
+
mlp_hidden: int = 1536
|
| 39 |
+
chunk: int = 256
|
| 40 |
+
seq_len: int = 512
|
| 41 |
+
tau_mem: float = 0.1
|
| 42 |
+
value_cap: float = 8.0
|
| 43 |
+
|
| 44 |
+
|
| 45 |
+
def _xavier(m):
|
| 46 |
+
if isinstance(m, nn.Linear):
|
| 47 |
+
bound = math.sqrt(6 / (m.in_features + m.out_features))
|
| 48 |
+
nn.init.uniform_(m.weight, -bound, bound)
|
| 49 |
+
if m.bias is not None:
|
| 50 |
+
nn.init.zeros_(m.bias)
|
| 51 |
+
|
| 52 |
+
|
| 53 |
+
def _poincare(x):
|
| 54 |
+
"""q -> q / (1 + |q|): maps R^n into the open unit ball."""
|
| 55 |
+
return x / (1.0 + x.norm(dim=-1, keepdim=True))
|
| 56 |
+
|
| 57 |
+
|
| 58 |
+
class PSSABlock(nn.Module):
|
| 59 |
+
"""SSM scan -> hyperbolic memory read -> gate -> plastic adapter -> MLP."""
|
| 60 |
+
|
| 61 |
+
def __init__(self, cfg):
|
| 62 |
+
super().__init__()
|
| 63 |
+
d, ds, dk, r = cfg.d_model, cfg.d_state, cfg.d_key, cfg.adapter_rank
|
| 64 |
+
self.cfg = cfg
|
| 65 |
+
self.norm = nn.Parameter(torch.ones(d))
|
| 66 |
+
self.norm_b = nn.Parameter(torch.zeros(d))
|
| 67 |
+
self.w_delta = nn.Linear(d, d, bias=False)
|
| 68 |
+
self.w_b = nn.Linear(d, ds, bias=False)
|
| 69 |
+
self.w_c = nn.Linear(d, ds, bias=False)
|
| 70 |
+
ratio = torch.arange(ds).float() / (ds - 1)
|
| 71 |
+
tau = torch.exp(math.log(1.5) + ratio * (math.log(200.0) - math.log(1.5)))
|
| 72 |
+
self.a_raw = nn.Parameter(torch.log(torch.expm1(1.0 / tau)).unsqueeze(0).repeat(d, 1))
|
| 73 |
+
self.w_qx = nn.Linear(d, dk, bias=False)
|
| 74 |
+
self.w_qh = nn.Linear(d, dk, bias=False)
|
| 75 |
+
self.w_gate = nn.Linear(d, d, bias=False)
|
| 76 |
+
self.w_proj = nn.Linear(d, d, bias=False)
|
| 77 |
+
self.ad_down = nn.Linear(d, r, bias=False)
|
| 78 |
+
self.ad_up = nn.Linear(r, d, bias=False)
|
| 79 |
+
self.mlp1 = nn.Linear(d, cfg.mlp_hidden, bias=False)
|
| 80 |
+
self.mlp2 = nn.Linear(cfg.mlp_hidden, d, bias=False)
|
| 81 |
+
self.mk_raw = nn.Parameter(torch.randn(cfg.mem_slots, dk) * 0.05)
|
| 82 |
+
self.mv = nn.Parameter(torch.randn(cfg.mem_slots, d) * 0.02)
|
| 83 |
+
for m in (self.w_delta, self.w_b, self.w_c, self.w_qx, self.w_qh,
|
| 84 |
+
self.w_gate, self.ad_down, self.mlp1):
|
| 85 |
+
_xavier(m)
|
| 86 |
+
nn.init.zeros_(self.w_proj.weight)
|
| 87 |
+
nn.init.zeros_(self.ad_up.weight)
|
| 88 |
+
nn.init.zeros_(self.mlp2.weight)
|
| 89 |
+
|
| 90 |
+
def forward(self, x, h0):
|
| 91 |
+
cfg = self.cfg
|
| 92 |
+
xn = self.norm * (x * torch.rsqrt(x.pow(2).mean(-1, keepdim=True) + 1e-5)) + self.norm_b
|
| 93 |
+
delta = F.softplus(self.w_delta(xn).float())
|
| 94 |
+
bp = self.w_b(xn).float()
|
| 95 |
+
cp = self.w_c(xn).float()
|
| 96 |
+
A = -F.softplus(self.a_raw)
|
| 97 |
+
y, h = _scan(xn, delta, bp, cp, A, h0)
|
| 98 |
+
|
| 99 |
+
q = _poincare((self.w_qx(xn) + self.w_qh(y)).float())
|
| 100 |
+
mk = _poincare(self.mk_raw)
|
| 101 |
+
qsq = (q * q).sum(-1, keepdim=True)
|
| 102 |
+
mksq = (mk * mk).sum(-1)
|
| 103 |
+
sq = (qsq + mksq.view(1, 1, -1) - 2 * (q @ mk.T)).clamp_min(0)
|
| 104 |
+
den = ((1 - qsq) * (1 - mksq).view(1, 1, -1)).clamp_min(1e-12)
|
| 105 |
+
dist = 2 * torch.asinh(torch.sqrt(sq / den))
|
| 106 |
+
w = torch.softmax(-dist / cfg.tau_mem, -1)
|
| 107 |
+
mv = cfg.value_cap * self.mv / (cfg.value_cap + self.mv.norm(dim=-1, keepdim=True))
|
| 108 |
+
m = w @ mv
|
| 109 |
+
|
| 110 |
+
g = torch.sigmoid(self.w_gate(xn))
|
| 111 |
+
z = y / math.sqrt(cfg.d_state) + g * self.w_proj(m.to(x.dtype))
|
| 112 |
+
z = z + self.ad_up(F.silu(self.ad_down(xn)))
|
| 113 |
+
z = z + self.mlp2(F.silu(self.mlp1(z)))
|
| 114 |
+
return z, h
|
| 115 |
+
|
| 116 |
+
|
| 117 |
+
class PSSAModel(nn.Module):
|
| 118 |
+
def __init__(self, cfg):
|
| 119 |
+
super().__init__()
|
| 120 |
+
self.cfg = cfg
|
| 121 |
+
self.embed = nn.Embedding(cfg.vocab_size, cfg.d_model)
|
| 122 |
+
self.blocks = nn.ModuleList([PSSABlock(cfg) for _ in range(cfg.depth)])
|
| 123 |
+
self.head = nn.Linear(cfg.d_model, cfg.vocab_size, bias=False)
|
| 124 |
+
self.scale = 1.0 / math.sqrt(cfg.depth)
|
| 125 |
+
|
| 126 |
+
def init_state(self, batch_size, device=None):
|
| 127 |
+
dev = device or self.embed.weight.device
|
| 128 |
+
return [torch.zeros(batch_size, self.cfg.d_model, self.cfg.d_state, device=dev)
|
| 129 |
+
for _ in range(self.cfg.depth)]
|
| 130 |
+
|
| 131 |
+
def forward_step(self, tokens, states):
|
| 132 |
+
x = self.embed(tokens)
|
| 133 |
+
for i, blk in enumerate(self.blocks):
|
| 134 |
+
z, h = blk(x, states[i])
|
| 135 |
+
x = z if i == 0 else x + z * self.scale
|
| 136 |
+
states[i] = h
|
| 137 |
+
return self.head(x) / math.sqrt(self.cfg.d_model), states
|
| 138 |
+
|
| 139 |
+
def forward(self, tokens):
|
| 140 |
+
states = self.init_state(tokens.shape[0], tokens.device)
|
| 141 |
+
outs = []
|
| 142 |
+
for s in range(0, tokens.shape[1], self.cfg.chunk):
|
| 143 |
+
logits, states = self.forward_step(tokens[:, s:s + self.cfg.chunk], states)
|
| 144 |
+
outs.append(logits)
|
| 145 |
+
return torch.cat(outs, 1)
|
| 146 |
+
|
| 147 |
+
|
| 148 |
+
def from_pretrained(model_dir: str, device: str = "cpu"):
|
| 149 |
+
"""Load the Hub export: (model, tokenizer). Works on CPU with no Triton install."""
|
| 150 |
+
from safetensors.torch import load_file
|
| 151 |
+
from tokenizers import Tokenizer
|
| 152 |
+
|
| 153 |
+
cfg = PSSAConfig()
|
| 154 |
+
sd = load_file(f"{model_dir}/model.safetensors")
|
| 155 |
+
sd = {k: v.float() for k, v in sd.items()}
|
| 156 |
+
m = PSSAModel(cfg)
|
| 157 |
+
m.load_state_dict(sd)
|
| 158 |
+
m.eval().to(device)
|
| 159 |
+
return m, Tokenizer.from_file(f"{model_dir}/tokenizer.json")
|
| 160 |
+
|
| 161 |
+
|
| 162 |
+
@torch.no_grad()
|
| 163 |
+
def generate(model, tokenizer, prompt, max_new=100, temp=0.7, top_k=40, eot=0, device="cpu"):
|
| 164 |
+
ids = tokenizer.encode(prompt).ids
|
| 165 |
+
states = model.init_state(1, device)
|
| 166 |
+
logits, states = model.forward_step(torch.tensor([ids], device=device), states)
|
| 167 |
+
out = []
|
| 168 |
+
for _ in range(max_new):
|
| 169 |
+
lg = logits[:, -1, :].float() / temp
|
| 170 |
+
if top_k:
|
| 171 |
+
kth = torch.topk(lg, top_k).values[:, [-1]]
|
| 172 |
+
lg = lg.masked_fill(lg < kth, -float("inf"))
|
| 173 |
+
nxt = torch.multinomial(torch.softmax(lg, -1), 1)
|
| 174 |
+
if int(nxt) == eot:
|
| 175 |
+
break
|
| 176 |
+
out.append(int(nxt))
|
| 177 |
+
logits, states = model.forward_step(nxt, states)
|
| 178 |
+
text = tokenizer.decode([i for i in out if i != eot])
|
| 179 |
+
return prompt + text
|
tokenizer.json
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|