mini-picoLM, in TypeScript
Rebuild the engine a third time: one ArrayBuffer, an arena, and eighteen steps to a browser that generates text with no server, no WASM and no dependencies. The hard part in JavaScript is the memory, not the maths.
- ArrayBuffer & arena
- typed-array views
- fused quant matmul
- async generation
- WebGPU
- shipping to a tab
Motivation
The problem
S21 proved the engine could be rebuilt in Python. Python is a comfortable place to be right: NumPy hides the memory and correctness is easy to check. It is also a place you cannot ship from, because nobody runs a Python process to render a web page.
So build it a third time, in TypeScript, with one constraint that changes everything: it has to run in a browser tab. No server, no WASM toolchain, no dependencies. One ArrayBuffer and eighteen steps.
This chapter prints every line of it — the whole engine, file by file, in the order you would write them. You can build the thing by following along and never open the repository. At the end there is a terminal, running the finished engine in this page, so you can check your own copy against it.
- steps, in dependency order
- 18
- ArrayBuffer → arena → … → WebGPU
- lines, all printed below
- 4018
- generated from the files that actually run
- runtime dependencies
- 0
- no WASM, no npm packages, no server
The source here cannot drift
src/lib/minipicolm/ by scripts/embed-minipicolm-source.mjs, and --check fails the build if the two disagree. This repository already learned that lesson from the hand-maintained line counts in chapters.ts. A chapter that prints several thousand lines of code needed a stronger guarantee than good intentions.Core idea
The build order, and the setup
mkdir -p minipicolm/{memory,tensor,model,runtime,tokenizer,format}
cd minipicolm
# there is nothing to install — that is the point
node --version # 22.6+ for --experimental-strip-typesTwo compiler settings matter. allowImportingTsExtensions lets relative imports keep their .ts suffix, which Node's ESM resolver requires; noEmit is its precondition. Together they let the same files run in a bundler and under node --experimental-strip-types with no build step.
{
"compilerOptions": {
"strict": true,
"noEmit": true,
"allowImportingTsExtensions": true,
"target": "ES2020",
"module": "esnext",
"moduleResolution": "bundler",
"lib": ["dom", "esnext"]
}
}Type-stripping erases, it does not compile
enum, namespace. Write constructor(private readonly dim: number) and the headless runner fails with ERR_UNSUPPORTED_TYPESCRIPT_SYNTAX. Declare the field and assign it, the way every class below does.Here is where you will end up:
src/lib/minipicolm/
├── memory/ 1-3 ArrayBuffer → arena → tensor views
│ ├── buffer.ts the single allocation
│ ├── arena.ts bump allocator, mark/release
│ └── tensor.ts shape + a view, nothing more
├── tensor/ 4,5,9,18
│ ├── matmul.ts matmul + the FUSED quantised path
│ ├── rmsnorm.ts x / rms(x) * w
│ ├── softmax.ts plain, and the streaming form
│ ├── activation.ts SwiGLU
│ └── webgpu.ts a real compute shader, feature-detected
├── model/ 6,7,10
│ ├── rope.ts interleaved rotation + precomputed tables
│ ├── attention.ts GQA over the cache, online softmax
│ ├── embedding.ts one dequantised row
│ ├── mlp.ts SwiGLU feed-forward
│ ├── block.ts the residual pair
│ └── llama.ts forward(token, pos)
├── runtime/ 8,12,13,17
│ ├── kv-cache.ts [layer][pos][kv_dim], written once
│ ├── workspace.ts every scratch buffer, allocated once
│ ├── sampler.ts temperature / top-k / top-p, seeded
│ ├── grammar.ts JSON pushdown automaton
│ └── generation.ts the async loop
├── tokenizer/ 11
│ └── bpe.ts SentencePiece, byte fallback
├── format/ 14-16
│ ├── gguf.ts read and write the container
│ └── quant.ts Q4_0, Q8_0, Q4_K, Q6_K
├── main.ts loadFromGGUF() and createToyModel()
├── shell.ts the commands behind the terminal
├── selftest.ts 90 assertions
├── run-selftest.mts the headless engine gate
└── run-shell-test.mts the headless shell gateMechanics
Steps 1–3 — memory
An inference engine's allocation pattern is completely known before it starts. You know the model shape and the context length, so you know every buffer you will ever need. That makes the simplest possible allocator the correct one.
memory/buffer.ts
The single allocation, plus alignment helpers. It fails loudly up front rather than at some random allocation later.
/**
* mini-picoLM — step 1: the ArrayBuffer
* =====================================
*
* Everything the engine touches lives in exactly one `ArrayBuffer`. Not one per
* tensor, not one per layer — one, allocated at startup and never grown.
*
* The reason is the same one picoLM has in C, and it survives the move to
* JavaScript intact: an inference engine's allocation pattern is completely
* known before the first token. You know the model shape, you know the context
* length, so you know every buffer you will ever need. Allocating them up front
* turns a stream of GC-visible objects into a single flat region.
*
* In the browser that matters more than it does in C. A `Float32Array` created
* per matmul is garbage the collector has to trace; a view into a pre-allocated
* buffer is not. The difference shows up as jitter between tokens, which is the
* one thing a streaming UI cannot hide.
*/
/** Typed-array constructors the arena knows how to hand out. */
export type TypedArrayCtor =
| Float32Array0
| Int32Array0
| Uint32Array0
| Uint8Array0
| Int8Array0
| Uint16Array0;
type Float32Array0 = typeof Float32Array;
type Int32Array0 = typeof Int32Array;
type Uint32Array0 = typeof Uint32Array;
type Uint8Array0 = typeof Uint8Array;
type Int8Array0 = typeof Int8Array;
type Uint16Array0 = typeof Uint16Array;
/** GPUs and SIMD both want 256-byte alignment; CPUs are happy with far less. */
export const ALIGNMENT = 256;
export function alignUp(bytes: number, alignment: number = ALIGNMENT): number {
return Math.ceil(bytes / alignment) * alignment;
}
export function bytesHuman(n: number): string {
if (n < 1024) return `${n} B`;
if (n < 1024 ** 2) return `${(n / 1024).toFixed(1)} KB`;
if (n < 1024 ** 3) return `${(n / 1024 ** 2).toFixed(1)} MB`;
return `${(n / 1024 ** 3).toFixed(2)} GB`;
}
/**
* Allocate the one buffer.
*
* Browsers cap a single ArrayBuffer well below the address space — the limit is
* around 2 GB on 64-bit Chrome and lower elsewhere — so this fails loudly
* rather than at some random allocation later.
*/
export function createBackingBuffer(bytes: number): ArrayBuffer {
if (!Number.isFinite(bytes) || bytes <= 0) {
throw new RangeError(`backing buffer size must be positive, got ${bytes}`);
}
try {
return new ArrayBuffer(bytes);
} catch (err) {
throw new RangeError(
`could not allocate ${bytesHuman(bytes)} in one ArrayBuffer. ` +
`Lower the context length or use a smaller model. (${String(err)})`,
);
}
}memory/arena.ts
A bump allocator. Allocation is one addition; there is no free list, no fragmentation, and no free() by design.
/**
* mini-picoLM — step 2: the arena allocator
* =========================================
*
* A bump allocator over the single ArrayBuffer from step 1. Allocation is one
* addition; there is no free list, no fragmentation, and no per-object header.
*
* There is also no `free()`. That sounds like a limitation and is actually the
* design: everything the engine allocates lives exactly as long as the engine
* does. Scratch space is reused by overwriting, not by releasing. The only
* lifetime operation that makes sense is `mark()` / `release()` — drop
* everything allocated after a point — which is how a request-scoped scratch
* region works.
*
* This is the same argument S06 makes about the KV cache from the other end.
* There, allocation had to be dynamic because output length is unknown. Here,
* every buffer's size is known before the first token, so the simplest possible
* allocator is also the right one.
*/
import { ALIGNMENT, alignUp, bytesHuman, createBackingBuffer } from "./buffer.ts";
export type Allocation = {
readonly name: string;
readonly byteOffset: number;
readonly byteLength: number;
};
export class Arena {
readonly buffer: ArrayBuffer;
readonly capacity: number;
private offset = 0;
private readonly log: Allocation[] = [];
constructor(capacityBytes: number) {
this.capacity = alignUp(capacityBytes);
this.buffer = createBackingBuffer(this.capacity);
}
/** Bytes handed out so far, including alignment padding. */
get used(): number {
return this.offset;
}
get remaining(): number {
return this.capacity - this.offset;
}
/** Every allocation, in order — what the S22 simulator draws. */
get allocations(): readonly Allocation[] {
return this.log;
}
/**
* Carve out `count` elements of `Ctor` and return a view onto them.
*
* The view aliases the arena; it is not a copy. Two views can legitimately
* overlap, which is how the workspace reuses one scratch region for several
* differently-shaped intermediates.
*/
alloc<T extends ArrayBufferView>(
Ctor: {
new (buffer: ArrayBuffer, byteOffset: number, length: number): T;
BYTES_PER_ELEMENT: number;
},
count: number,
name = "anon",
): T {
const byteLength = count * Ctor.BYTES_PER_ELEMENT;
const byteOffset = alignUp(this.offset, ALIGNMENT);
if (byteOffset + byteLength > this.capacity) {
throw new RangeError(
`arena exhausted allocating ${bytesHuman(byteLength)} for "${name}": ` +
`${bytesHuman(this.remaining)} left of ${bytesHuman(this.capacity)}`,
);
}
this.offset = byteOffset + byteLength;
this.log.push({ name, byteOffset, byteLength });
return new Ctor(this.buffer, byteOffset, count);
}
f32(count: number, name?: string): Float32Array {
return this.alloc(Float32Array, count, name);
}
i32(count: number, name?: string): Int32Array {
return this.alloc(Int32Array, count, name);
}
u8(count: number, name?: string): Uint8Array {
return this.alloc(Uint8Array, count, name);
}
/** Remember the current watermark. */
mark(): number {
return this.offset;
}
/**
* Drop everything allocated since `mark`. The bytes are not cleared — the
* next allocation simply overwrites them, which is the whole point.
*/
release(mark: number): void {
if (mark < 0 || mark > this.offset) {
throw new RangeError(`bad arena mark ${mark} (offset is ${this.offset})`);
}
this.offset = mark;
while (this.log.length && this.log[this.log.length - 1].byteOffset >= mark) {
this.log.pop();
}
}
summary(): string {
return `${bytesHuman(this.used)} / ${bytesHuman(this.capacity)} in ${this.log.length} allocations`;
}
}memory/tensor.ts
A shape and a view. Inference needs no autograd, no broadcasting and no lazy graph, and each of them would cost an allocation per operation.
/**
* mini-picoLM — step 3: tensor views
* ==================================
*
* A tensor here is a shape plus a `Float32Array` that already points into the
* arena. It owns nothing and allocates nothing.
*
* Note what is deliberately missing: no autograd, no broadcasting, no dtype
* promotion, no lazy graph. Inference needs none of it, and every one of those
* features costs an allocation or an indirection per operation. The whole type
* is twenty lines because the whole job is "interpret these bytes as a matrix".
*/
import type { Arena } from "./arena.ts";
export type Shape = readonly number[];
export class Tensor {
readonly data: Float32Array;
readonly shape: Shape;
readonly name: string;
constructor(data: Float32Array, shape: Shape, name = "anon") {
const n = shape.reduce((a, b) => a * b, 1);
if (data.length !== n) {
throw new RangeError(
`tensor "${name}": shape [${shape}] needs ${n} elements, view has ${data.length}`,
);
}
this.data = data;
this.shape = shape;
this.name = name;
}
static zeros(arena: Arena, shape: Shape, name = "anon"): Tensor {
const n = shape.reduce((a, b) => a * b, 1);
return new Tensor(arena.f32(n, name), shape, name);
}
/** Wrap an existing view without copying — the common case. */
static of(data: Float32Array, shape: Shape, name = "anon"): Tensor {
return new Tensor(data, shape, name);
}
get rank(): number {
return this.shape.length;
}
get length(): number {
return this.data.length;
}
get rows(): number {
return this.rank >= 2 ? this.shape[0] : 1;
}
get cols(): number {
return this.shape[this.rank - 1];
}
/** Row `i` of a 2-D tensor, as a view. Still no copy. */
row(i: number): Float32Array {
const c = this.cols;
return this.data.subarray(i * c, (i + 1) * c);
}
/** A sub-range of the last axis, as a view — used to slice heads apart. */
slice(start: number, end: number): Float32Array {
return this.data.subarray(start, end);
}
fill(v: number): this {
this.data.fill(v);
return this;
}
copyFrom(src: ArrayLike<number>): this {
this.data.set(src);
return this;
}
toString(): string {
return `Tensor(${this.name}, [${this.shape}])`;
}
}The arena is live in this page. mem lists every allocation it has handed out, in order: the KV cache first, then each workspace buffer.
The point is the collector, not the speed
malloc calls. In JavaScript it saves something more valuable. A Float32Array created per matmul is garbage the collector must trace, and a collection between two tokens is a visible stall in a streaming UI. Because every view aliases one buffer allocated up front, the token loop allocates nothing at all, and inter-token latency stays flat rather than merely low.Explore
Steps 4, 5, 9 — the kernels
Four small files. matmul carries the whole performance story; the rest are one-liners with a comment explaining why they are shaped the way they are.
tensor/matmul.ts
Matrix-vector, because decode has no batch dimension. Row-major and sequential so the prefetcher keeps up, four-way unrolled so the adds are independent, and never allocating: the output view is passed in.
/**
* mini-picoLM — step 4: matmul
* ============================
*
* Decode is one token at a time, so every matmul in this engine is really a
* matrix-vector product: `out[d] = dot(W[d], x)`. There is no batch dimension,
* exactly as in picoLM, and for the same reason — one user, one core.
*
* That makes this function completely memory-bound, which is S05's whole point
* arriving in JavaScript. The arithmetic is trivial; the cost is streaming `W`
* through the CPU once per token. Two consequences shape the code below:
*
* * **Row-major, sequential.** `W` is walked in the order it is stored, so
* the prefetcher can keep up. Transposing the loops costs ~5× on the same
* arithmetic.
* * **No allocation.** The output view is passed in. A function that returned
* a fresh Float32Array would allocate `n_layers × 7` arrays per token.
*
* The quantised variants dequantise inline and never materialise a float copy
* of `W` — the fused path S08 argues for and the Python port of S21 could not
* express. TypeScript can, because a hand-written loop over a Uint8Array is
* exactly as fast as any other loop here.
*/
import { dequantBlockInto, blockSize, type QuantType } from "../format/quant.ts";
/** out[d] = dot(W[d, :], x), W row-major [d, n], float32. */
export function matmul(
out: Float32Array,
x: Float32Array,
w: Float32Array,
n: number,
d: number,
): void {
for (let i = 0; i < d; i++) {
const base = i * n;
let acc = 0;
// Four-way unrolled: JS engines do not vectorise this for us, but they do
// keep four independent accumulators in registers, which hides the latency
// of the dependent add chain.
let j = 0;
let a0 = 0;
let a1 = 0;
let a2 = 0;
let a3 = 0;
for (; j + 3 < n; j += 4) {
a0 += w[base + j] * x[j];
a1 += w[base + j + 1] * x[j + 1];
a2 += w[base + j + 2] * x[j + 2];
a3 += w[base + j + 3] * x[j + 3];
}
acc = a0 + a1 + a2 + a3;
for (; j < n; j++) acc += w[base + j] * x[j];
out[i] = acc;
}
}
/**
* The fused path: `W` is still quantised bytes, and one block is decoded into a
* small stack scratch buffer immediately before it is consumed.
*
* Nothing full-precision the size of `W` is ever stored, so the bytes crossing
* the memory bus are the quantised ones — which is the entire reason to
* quantise. Compare `dequantizeAll` + `matmul`, which reads the int4 weights
* AND writes and re-reads a float32 copy: strictly more traffic than never
* quantising at all.
*/
export function matmulQuant(
out: Float32Array,
x: Float32Array,
w: Uint8Array,
n: number,
d: number,
qtype: QuantType,
scratch: Float32Array,
): void {
const bs = blockSize(qtype);
if (n % bs !== 0) {
throw new RangeError(`matmulQuant: ${n} is not a multiple of the ${bs}-weight block`);
}
const blocksPerRow = n / bs;
for (let i = 0; i < d; i++) {
let acc = 0;
for (let b = 0; b < blocksPerRow; b++) {
// decode 32 or 256 weights into registers-worth of scratch, use, discard
dequantBlockInto(w, i * blocksPerRow + b, qtype, scratch);
const off = b * bs;
for (let j = 0; j < bs; j++) acc += scratch[j] * x[off + j];
}
out[i] = acc;
}
}
/** out += a (in place) — the residual stream's only operation. */
export function addInto(out: Float32Array, a: Float32Array): void {
for (let i = 0; i < out.length; i++) out[i] += a[i];
}
/** out[i] = a[i] * b[i] — SwiGLU's gate. */
export function mulInto(out: Float32Array, a: Float32Array, b: Float32Array): void {
for (let i = 0; i < out.length; i++) out[i] = a[i] * b[i];
}
export function dot(a: Float32Array, b: Float32Array, n: number): number {
let acc = 0;
for (let i = 0; i < n; i++) acc += a[i] * b[i];
return acc;
}tensor/rmsnorm.ts
LayerNorm minus the parts that turned out not to matter. Note the epsilon. It is model config shipped in GGUF metadata, and the values differ: TinyLlama and Llama-2 train with 1e-5, the original LLaMA with S03's 1e-6. Load with the wrong one and every activation shifts slightly, compounding over the layers.
/**
* mini-picoLM — step 5: RMSNorm
* =============================
*
* LayerNorm with the parts that turned out not to matter removed: no mean
* subtraction, no bias. `x / rms(x) * weight`.
*
* The epsilon is inside the square root and is not decorative. It is model
* configuration, shipped in GGUF metadata: TinyLlama and Llama-2 train with
* 1e-5, the original LLaMA with 1e-6 (S03's value). Load a model with the
* wrong one and every activation shifts slightly, compounding over 22 layers —
* the kind of mismatch that produces a model which works and is quietly a
* little worse.
*/
export const RMS_EPS = 1e-5;
export function rmsnorm(
out: Float32Array,
x: Float32Array,
weight: Float32Array,
eps: number = RMS_EPS,
): void {
const n = x.length;
let ss = 0;
for (let i = 0; i < n; i++) ss += x[i] * x[i];
const scale = 1 / Math.sqrt(ss / n + eps);
for (let i = 0; i < n; i++) out[i] = x[i] * scale * weight[i];
}tensor/softmax.ts
Two shapes of the same function. OnlineSoftmax is the one that matters. It never sees the whole row, so attention can stream over a cache of any length without allocating.
/**
* mini-picoLM — softmax, twice
* ============================
*
* Two shapes of the same function, and the difference is the subject of S12.
*
* `softmax` is the ordinary one: find the max, exponentiate, normalise. It
* needs the whole row in memory at once.
*
* `OnlineSoftmax` never sees the whole row. It keeps a running maximum and a
* running sum, and rescales both whenever a larger value arrives. That is what
* lets attention stream over a KV cache of any length without ever allocating
* an array of scores — the single most important structural idea in a decode
* kernel, and it is nine lines.
*/
/** In-place softmax over `x[0..n)`. Subtracting the max is what stops exp() overflowing. */
export function softmax(x: Float32Array, n: number = x.length): void {
let max = -Infinity;
for (let i = 0; i < n; i++) if (x[i] > max) max = x[i];
let sum = 0;
for (let i = 0; i < n; i++) {
const e = Math.exp(x[i] - max);
x[i] = e;
sum += e;
}
const inv = 1 / sum;
for (let i = 0; i < n; i++) x[i] *= inv;
}
/**
* The streaming form. Feed it one score and one value vector at a time; the
* accumulator holds the correctly-weighted partial result at every point.
*
* m running maximum
* l running sum of exp(score - m)
* acc running sum of exp(score - m) * v
*
* When a score exceeds `m`, everything accumulated so far was scaled against
* the old maximum, so it is corrected by exp(m_old - m_new) — one multiply per
* element, and the invariant holds again.
*/
export class OnlineSoftmax {
private m = -Infinity;
private l = 0;
private readonly acc: Float32Array;
private readonly dim: number;
constructor(dim: number, scratch: Float32Array) {
this.dim = dim;
if (scratch.length < dim) {
throw new RangeError(`OnlineSoftmax needs ${dim} scratch elements, got ${scratch.length}`);
}
this.acc = scratch.subarray(0, dim);
}
reset(): void {
this.m = -Infinity;
this.l = 0;
this.acc.fill(0);
}
/** Fold one (score, value) pair into the accumulator. */
push(score: number, v: Float32Array, vOffset: number): void {
// A -Infinity score is a masked position and its weight is exp(-inf) = 0,
// so drop it at the door. Without this, a fully-masked stream leaves m at
// -Infinity and the else-branch below computes exp(-inf - -inf) = NaN —
// the classic fully-masked-tile bug. Finite scores need no guard at all:
// exp(-inf - score) is already 0.
if (score === -Infinity) return;
const { acc, dim } = this;
if (score > this.m) {
const correction = Math.exp(this.m - score);
this.l = this.l * correction + 1;
for (let i = 0; i < dim; i++) acc[i] = acc[i] * correction + v[vOffset + i];
this.m = score;
} else {
const wgt = Math.exp(score - this.m);
this.l += wgt;
for (let i = 0; i < dim; i++) acc[i] += wgt * v[vOffset + i];
}
}
/** Normalise into `out`. Safe to call once per query head. */
finish(out: Float32Array, outOffset: number): void {
const inv = this.l === 0 ? 0 : 1 / this.l;
for (let i = 0; i < this.dim; i++) out[outOffset + i] = this.acc[i] * inv;
}
}tensor/activation.ts
SwiGLU. Three matrices where a classic FFN has two; the extra one is a gate, so the network can suppress a channel rather than merely fail to excite it.
/**
* mini-picoLM — step 9: SwiGLU
* ============================
*
* `SwiGLU(x) = (SiLU(x·W1) * (x·W3)) · W2`
*
* Three matrices where a classic FFN has two. The extra one is the *gate*: its
* output multiplies the activation elementwise, so the network can suppress a
* channel rather than merely fail to excite it.
*
* The cost is the reason `d_ff` looks odd in Llama configs. A two-matrix FFN at
* `4 * d_model` and a three-matrix one at `(8/3) * d_model` hold the same
* parameter count, which is why you see 5632 next to 2048 rather than 8192.
*/
/** SiLU, a.k.a. swish: x * sigmoid(x). */
export function silu(x: Float32Array, n: number = x.length): void {
for (let i = 0; i < n; i++) x[i] = x[i] / (1 + Math.exp(-x[i]));
}
export function siluScalar(v: number): number {
return v / (1 + Math.exp(-v));
}
/**
* Fuse the activation and the gate into one pass.
*
* `hb` holds the gate projection, `hb2` the up projection. Doing this in one
* loop rather than two keeps `hb` in cache between the two reads — the same
* argument as kernel fusion in S13, at a scale where you can see the whole
* thing.
*/
export function siluMulInto(hb: Float32Array, hb2: Float32Array, n: number = hb.length): void {
for (let i = 0; i < n; i++) {
const g = hb[i];
hb[i] = (g / (1 + Math.exp(-g))) * hb2[i];
}
}bench runs the kernels above for real: sixteen forward passes, timed, with the matmul and MAC counts they cost.
The masked position is the special case, not the first fold
exp(-Infinity - score) is 0 all by itself, so the first fold needs no guard. The case that does need one is a fully-masked stream: every score -Infinity, the running max never moves, and the running-sum branch computes exp(-Infinity - -Infinity), which is NaN. Real kernels guard exactly this, and so does push: a -Infinity score is dropped at the door, because its weight is exp(-inf) = 0 anyway.Build it
Steps 15–16 — Q8 and Q4, early
These are steps 15 and 16 in the numbering, but the matmul above already imports them, so they belong here. The formats are the ones from S21, compiled against picoLM's own quant.c and compared byte for byte on random blocks.
format/quant.ts
Q4_0 and Q8_0 are 32 weights and one fp16 scale. Q4_K and Q6_K are 256-weight super-blocks with nested scales: Q4_K a 6-bit scale and min per 32-weight sub-block, Q6_K sixteen signed 8-bit scales, one per 16-weight sub-block. That nesting is what tracks outliers at four bits, and why a real 4-bit model is 4.5 bits per weight.
/**
* mini-picoLM — steps 15 & 16: Q8 and Q4
* ======================================
*
* The GGUF block formats, in TypeScript. These are the same algorithms the S21
* Python port implements, and that port was compiled against picoLM's own
* `quant.c` and compared byte for byte on random blocks — so the numbers here
* have a verified reference. `selftest.ts` re-checks them against blocks
* written by hand from the spec.
*
* Two families:
*
* **Q8_0 / Q4_0** — 32 weights, one fp16 scale. Simple, symmetric, and what
* "4-bit" meant before 2023.
*
* **Q4_K / Q6_K** — 256-weight super-blocks with *nested* scales. Q4_K packs
* a 6-bit scale and min per 32-weight sub-block under one fp16 pair; Q6_K
* carries sixteen signed 8-bit scales, one per 16-weight sub-block, under a
* single fp16 scale. That nesting is what tracks outliers at four bits, and
* it is why a real 4-bit model is 4.5 bits per weight rather than 4.0.
*
* The awkward bit-packing in `getScaleMin` is the format, not the port: for the
* first four sub-blocks the 6-bit fields sit in their own bytes, and for the
* last four they straddle byte boundaries.
*/
export type QuantType = "F32" | "F16" | "Q4_0" | "Q8_0" | "Q4_K" | "Q6_K";
/** GGUF's numeric tensor-type tags. */
export const GGUF_TYPE: Record<number, QuantType> = {
0: "F32",
1: "F16",
2: "Q4_0",
8: "Q8_0",
12: "Q4_K",
14: "Q6_K",
};
const BLOCK: Record<QuantType, number> = {
F32: 1,
F16: 1,
Q4_0: 32,
Q8_0: 32,
Q4_K: 256,
Q6_K: 256,
};
const QSIZE: Record<QuantType, number> = {
F32: 4,
F16: 2,
Q4_0: 18,
Q8_0: 34,
Q4_K: 144,
Q6_K: 210,
};
export function blockSize(t: QuantType): number {
return BLOCK[t];
}
export function quantSize(t: QuantType): number {
return QSIZE[t];
}
export function rowSize(t: QuantType, n: number): number {
return (n / BLOCK[t]) * QSIZE[t];
}
export function bitsPerWeight(t: QuantType): number {
return (8 * QSIZE[t]) / BLOCK[t];
}
/* ------------------------------------------------------------------ fp16 */
/**
* JavaScript has no Float16Array in every engine we care about yet, so decode
* IEEE-754 binary16 by hand — the same bit surgery picoLM writes in C because
* its target CPUs lack hardware fp16.
*/
export function fp16ToFp32(h: number): number {
const sign = (h >> 15) & 1;
const exp = (h >> 10) & 0x1f;
const mant = h & 0x3ff;
let value: number;
if (exp === 0) {
value = mant === 0 ? 0 : mant * 2 ** -24; // subnormal
} else if (exp === 31) {
value = mant === 0 ? Infinity : NaN;
} else {
value = (mant + 1024) * 2 ** (exp - 25);
}
return sign ? -value : value;
}
const _f32 = new Float32Array(1);
const _u32 = new Uint32Array(_f32.buffer);
export function fp32ToFp16(f: number): number {
_f32[0] = f;
const bits = _u32[0];
const sign = (bits >>> 16) & 0x8000;
const exp = ((bits >>> 23) & 0xff) - 127 + 15;
const mant = bits & 0x7fffff;
if (((bits >>> 23) & 0xff) === 0xff) return sign | 0x7c00 | (mant ? 0x200 : 0);
if (exp >= 31) return sign | 0x7c00;
if (exp <= 0) {
if (exp < -10) return sign;
const m = (mant | 0x800000) >>> (14 - exp);
return sign | m;
}
return sign | (exp << 10) | (mant >>> 13);
}
/* ------------------------------------------------- one block at a time */
/** Q4_K's six-bit scale/min pair for sub-block `j`, unpacked from 12 bytes. */
function getScaleMin(q: Uint8Array, base: number, j: number): [number, number] {
if (j < 4) {
return [q[base + j] & 63, q[base + j + 4] & 63];
}
const sc = (q[base + j + 4] & 0xf) | ((q[base + j - 4] >> 6) << 4);
const mn = (q[base + j + 4] >> 4) | ((q[base + j] >> 6) << 4);
return [sc, mn];
}
/**
* Decode block `index` of a quantised tensor into `out[0..blockSize)`.
*
* This is the primitive the fused matmul calls. It writes into a caller-owned
* scratch buffer so that nothing the size of the weight matrix is ever
* allocated — the whole point of fusing.
*/
export function dequantBlockInto(
raw: Uint8Array,
index: number,
qtype: QuantType,
out: Float32Array,
): void {
const off = index * QSIZE[qtype];
const dv = new DataView(raw.buffer, raw.byteOffset, raw.byteLength);
switch (qtype) {
case "F32": {
out[0] = dv.getFloat32(off, true);
return;
}
case "F16": {
out[0] = fp16ToFp32(dv.getUint16(off, true));
return;
}
case "Q8_0": {
const d = fp16ToFp32(dv.getUint16(off, true));
for (let i = 0; i < 32; i++) out[i] = d * dv.getInt8(off + 2 + i);
return;
}
case "Q4_0": {
const d = fp16ToFp32(dv.getUint16(off, true));
for (let i = 0; i < 16; i++) {
const b = raw[off + 2 + i];
out[i] = d * ((b & 0xf) - 8);
out[i + 16] = d * ((b >> 4) - 8);
}
return;
}
case "Q4_K": {
const d = fp16ToFp32(dv.getUint16(off, true));
const dmin = fp16ToFp32(dv.getUint16(off + 2, true));
const sBase = off + 4;
const qBase = off + 16;
let o = 0;
let qb = 0;
for (let j = 0; j < 4; j++) {
const [sc1, mn1] = getScaleMin(raw, sBase, j * 2);
const [sc2, mn2] = getScaleMin(raw, sBase, j * 2 + 1);
const d1 = d * sc1;
const m1 = dmin * mn1;
const d2 = d * sc2;
const m2 = dmin * mn2;
for (let l = 0; l < 32; l++) {
const b = raw[qBase + qb + l];
out[o + l] = d1 * (b & 0xf) - m1;
out[o + l + 32] = d2 * (b >> 4) - m2;
}
o += 64;
qb += 32;
}
return;
}
case "Q6_K": {
const d = fp16ToFp32(dv.getUint16(off + 208, true));
const qlBase = off;
const qhBase = off + 128;
const scBase = off + 192;
for (let c = 0; c < 2; c++) {
const ql = qlBase + c * 64;
const qh = qhBase + c * 32;
const isBase = c * 8;
const y = c * 128;
for (let l = 0; l < 32; l++) {
const isL = isBase + (l >> 4);
const lo = raw[ql + l];
const hi = raw[ql + l + 32];
const h = raw[qh + l];
out[y + l] = d * dv.getInt8(scBase + isL) * (((lo & 0xf) | (((h >> 0) & 3) << 4)) - 32);
out[y + l + 32] =
d * dv.getInt8(scBase + isL + 2) * (((hi & 0xf) | (((h >> 2) & 3) << 4)) - 32);
out[y + l + 64] =
d * dv.getInt8(scBase + isL + 4) * (((lo >> 4) | (((h >> 4) & 3) << 4)) - 32);
out[y + l + 96] =
d * dv.getInt8(scBase + isL + 6) * (((hi >> 4) | (((h >> 6) & 3) << 4)) - 32);
}
}
return;
}
default: {
const never: never = qtype;
throw new Error(`unsupported quant type ${String(never)}`);
}
}
}
/**
* Decode `n` consecutive weights.
*
* Convenient, and exactly what a hot loop must not do — it materialises the
* float copy that fusing exists to avoid. Used for the norm weights (tiny) and
* by the tests.
*/
export function dequantizeRow(raw: Uint8Array, n: number, qtype: QuantType): Float32Array {
const bs = BLOCK[qtype];
if (n % bs !== 0) {
throw new RangeError(`${n} weights is not a whole number of ${qtype} blocks`);
}
const out = new Float32Array(n);
const scratch = new Float32Array(bs);
const nb = n / bs;
for (let b = 0; b < nb; b++) {
dequantBlockInto(raw, b, qtype, scratch);
out.set(scratch, b * bs);
}
return out;
}Run quant for the block table, or quant Q4_K to decode an actual super-block. fp16 65504 shows the largest finite half; try fp16 1e-8 for the subnormal floor and fp16 70000 for saturation.
This is where TypeScript beats NumPy
Uint8Array is exactly as fast as any other loop here, so matmulQuant above simply writes the fused loop out. The browser port gets the win the Python port had to give up.Production notes
Steps 6–10 — the model
RoPE, the cache, attention, and the block. Six files that turn the kernels into a network.
model/rope.ts
Rotary position embedding, interleaved, pairing x[2i] with x[2i+1]. S03 teaches the NeoX convention, which pairs x[i] with x[i + D/2]. Both are called RoPE, both are self-consistent, and a model run with the wrong one produces fluent, confident, subtly wrong text.
/**
* mini-picoLM — step 6: RoPE
* ==========================
*
* Rotary position embedding. Instead of adding a position vector, rotate each
* pair of dimensions by an angle proportional to the position. The dot product
* of two rotated vectors then depends only on the *gap* between their
* positions — which is exactly why a KV cache is sound: a cached key's rotation
* is fixed forever, because its absolute position never changes.
*
* There are two incompatible conventions, both called RoPE:
*
* **interleaved** — pair `x[2i]` with `x[2i+1]`. llama.cpp, picoLM, and every
* Llama GGUF file.
*
* **NeoX / split-half** — pair `x[i]` with `x[i + D/2]`. GPT-NeoX, and the
* one S03 teaches.
*
* Pick the wrong one and the model runs, sounds fluent, and is subtly wrong.
* This is the interleaved one, because this engine loads GGUF.
*/
/** Precomputed cos/sin for every position — keeps trig out of the token loop. */
export class RopeTables {
readonly cos: Float32Array;
readonly sin: Float32Array;
readonly half: number;
constructor(maxSeq: number, headDim: number, freqBase = 10000) {
this.half = headDim >> 1;
this.cos = new Float32Array(maxSeq * this.half);
this.sin = new Float32Array(maxSeq * this.half);
for (let pos = 0; pos < maxSeq; pos++) {
for (let i = 0; i < this.half; i++) {
const theta = pos / Math.pow(freqBase, (2 * i) / headDim);
this.cos[pos * this.half + i] = Math.cos(theta);
this.sin[pos * this.half + i] = Math.sin(theta);
}
}
}
}
/**
* Rotate `x` in place: `nHeads` heads of `headDim` laid out contiguously.
*
* Applied to Q and K, never to V — V carries content, not position.
*/
export function ropeInPlace(
x: Float32Array,
nHeads: number,
headDim: number,
tables: RopeTables,
pos: number,
): void {
const half = tables.half;
const base = pos * half;
for (let h = 0; h < nHeads; h++) {
const o = h * headDim;
for (let i = 0; i < half; i++) {
const c = tables.cos[base + i];
const s = tables.sin[base + i];
const x0 = x[o + i * 2];
const x1 = x[o + i * 2 + 1];
x[o + i * 2] = x0 * c - x1 * s;
x[o + i * 2 + 1] = x0 * s + x1 * c;
}
}
}runtime/kv-cache.ts
One flat region indexed [layer][position][kv_dim]. Keys are stored after rotation, because a cached token's position never changes; entries are immutable once written.
/**
* mini-picoLM — step 8: the KV cache
* ==================================
*
* One flat region of the arena, indexed `[layer][position][kv_dim]`. Writing
* position `p` is a `set()` into a subarray; reading positions `0..p` is a
* subarray. No allocation on either path.
*
* Two properties do all the work:
*
* **Keys are stored after RoPE.** A cached token's position never changes, so
* its rotation never changes. Storing pre-rotation would force re-rotating
* the whole cache every step — replacing one compute cost with a smaller one
* and defeating most of the point.
*
* **Entries are immutable once written.** Writing position N must not disturb
* anything before it. `selftest.ts` asserts this, because a cache that
* silently corrupts history produces text that degrades over hundreds of
* tokens with no error anywhere.
*
* Unlike picoLM this stores float32, not fp16 — `Float16Array` is not yet
* available across the browsers this has to run in. That doubles the cache, and
* since the cache is the resident cost, it is the main reason the context
* length here is smaller than picoLM's on the same device. The arithmetic is in
* the S22 simulator.
*/
import type { Arena } from "../memory/arena.ts";
export class KVCache {
readonly key: Float32Array;
readonly value: Float32Array;
readonly nLayers: number;
readonly maxSeq: number;
readonly kvDim: number;
private _length = 0;
constructor(arena: Arena, nLayers: number, maxSeq: number, kvDim: number) {
this.nLayers = nLayers;
this.maxSeq = maxSeq;
this.kvDim = kvDim;
const n = nLayers * maxSeq * kvDim;
this.key = arena.f32(n, "kv.key");
this.value = arena.f32(n, "kv.value");
}
/** Positions written so far. */
get length(): number {
return this._length;
}
get bytes(): number {
return this.key.byteLength + this.value.byteLength;
}
private offset(layer: number, pos: number): number {
return (layer * this.maxSeq + pos) * this.kvDim;
}
/** Append this layer's K and V for position `pos`. */
write(layer: number, pos: number, k: Float32Array, v: Float32Array): void {
if (pos >= this.maxSeq) {
throw new RangeError(`position ${pos} past the ${this.maxSeq}-token cache`);
}
const o = this.offset(layer, pos);
this.key.set(k, o);
this.value.set(v, o);
if (layer === this.nLayers - 1) this._length = Math.max(this._length, pos + 1);
}
/** Byte offset of one head's K vector — attention reads through this. */
headOffset(layer: number, pos: number, kvHead: number, headDim: number): number {
return this.offset(layer, pos) + kvHead * headDim;
}
reset(): void {
this._length = 0;
this.key.fill(0);
this.value.fill(0);
}
}model/attention.ts
GQA over the cache with an online softmax. There is no score array, so peak intermediate memory is head_dim floats whatever the context length.
/**
* mini-picoLM — step 7: attention
* ===============================
*
* Grouped-query attention over the KV cache, computed with an online softmax.
*
* The shape of this function is the whole argument of S12. There is no `scores`
* array. At no context length does this allocate anything proportional to the
* sequence — it streams over the cache once per head, folding each (score,
* value) pair into a running accumulator. Peak intermediate memory is
* `head_dim` floats, whatever the context.
*
* GQA is the `kvHead = h / kvMul` line. Several query heads read the same K and
* V rows, so the cache shrinks by that factor while the query side keeps its
* full expressiveness. It is a memory optimisation, not a quality one.
*/
import { OnlineSoftmax } from "../tensor/softmax.ts";
import type { KVCache } from "../runtime/kv-cache.ts";
export type AttentionShape = {
nHeads: number;
nKvHeads: number;
headDim: number;
};
/**
* out[0..nHeads*headDim) = attention(q, cache[layer][0..=pos])
*
* `scratch` must hold at least `headDim` floats and is reused across heads.
*/
export function attention(
out: Float32Array,
q: Float32Array,
cache: KVCache,
layer: number,
pos: number,
shape: AttentionShape,
scratch: Float32Array,
): void {
const { nHeads, nKvHeads, headDim } = shape;
const kvMul = nHeads / nKvHeads;
const inv = 1 / Math.sqrt(headDim);
const online = new OnlineSoftmax(headDim, scratch);
for (let h = 0; h < nHeads; h++) {
const kvHead = Math.floor(h / kvMul);
const qo = h * headDim;
online.reset();
for (let t = 0; t <= pos; t++) {
const ko = cache.headOffset(layer, t, kvHead, headDim);
let score = 0;
for (let i = 0; i < headDim; i++) score += q[qo + i] * cache.key[ko + i];
online.push(score * inv, cache.value, ko);
}
online.finish(out, qo);
}
}model/embedding.ts
One dequantised row. The embedding table is usually the largest tensor in the file and a decode step touches exactly one row of it.
/**
* mini-picoLM — the embedding lookup
* ==================================
*
* A row of the embedding table, dequantised on demand.
*
* Worth noticing how little happens here. `token_embd.weight` is usually the
* largest single tensor in the file — 32000 × 2048 for TinyLlama — and a decode
* step touches exactly one of its rows. Dequantising the whole table at load
* time would cost hundreds of megabytes to save a few microseconds per token.
* So it stays quantised, mapped, and untouched except for the one row.
*/
import { dequantBlockInto, blockSize, rowSize, type QuantType } from "../format/quant.ts";
/**
* Write row `token` of a [vocab, dim] quantised table into `out`.
*
* `out` is a caller-owned view — the residual stream itself, in practice, so
* the embedding is written straight into the buffer the first layer reads.
*/
export function embedInto(
out: Float32Array,
table: Uint8Array,
token: number,
dim: number,
qtype: QuantType,
scratch: Float32Array,
): void {
const bs = blockSize(qtype);
const perRow = rowSize(qtype, dim);
const rowStart = token * perRow;
const blocksPerRow = dim / bs;
// `dequantBlockInto` indexes blocks from the start of the view it is given,
// so hand it a view that starts at this row.
const row = table.subarray(rowStart, rowStart + perRow);
for (let b = 0; b < blocksPerRow; b++) {
dequantBlockInto(row, b, qtype, scratch);
out.set(scratch.subarray(0, bs), b * bs);
}
}model/mlp.ts
Gate and up in parallel, activate-and-multiply, then down. Roughly 70% of a Llama block's parameters live here.
/**
* mini-picoLM — the feed-forward network
* ======================================
*
* SwiGLU: gate and up in parallel, activate-and-multiply, then down.
*
* This is where most of the model's parameters live — roughly 70% of a Llama
* block — which makes it the dominant term in decode bandwidth. Everything the
* quantisation chapters buy, they buy mostly here.
*/
import { siluMulInto } from "../tensor/activation.ts";
import type { QuantWeight } from "./llama.ts";
import type { Workspace } from "../runtime/workspace.ts";
export function feedForward(
out: Float32Array,
x: Float32Array,
gate: QuantWeight,
up: QuantWeight,
down: QuantWeight,
dim: number,
hidden: number,
ws: Workspace,
): void {
ws.matmul(ws.hb, x, gate, dim, hidden);
ws.matmul(ws.hb2, x, up, dim, hidden);
siluMulInto(ws.hb, ws.hb2, hidden);
ws.matmul(out, ws.hb, down, hidden, dim);
}runtime/workspace.ts
Every scratch buffer, allocated once, plus the single matmul entry point that dispatches on how a weight is stored.
/**
* mini-picoLM — the workspace
* ===========================
*
* Every scratch buffer the forward pass needs, allocated once from the arena.
*
* The list is short and fixed because the shapes are known before the first
* token: one residual stream, two same-size temporaries, a query vector, two
* FFN-width buffers, a logits row, and one block of dequantisation scratch.
* That is the entire working set of a decode step.
*
* `matmul` lives here rather than in `tensor/` because it needs the scratch
* buffer to fuse against. Routing every weight through this one method also
* makes the WebGPU backend a single substitution rather than a rewrite.
*/
import type { Arena } from "../memory/arena.ts";
import { matmul, matmulQuant } from "../tensor/matmul.ts";
import { blockSize } from "../format/quant.ts";
import type { QuantWeight } from "../model/llama.ts";
export class Workspace {
readonly x: Float32Array; // residual stream
readonly xb: Float32Array; // post-norm / attention out
readonly xb2: Float32Array; // second temporary
readonly q: Float32Array;
readonly k: Float32Array;
readonly v: Float32Array;
readonly hb: Float32Array; // ffn hidden (gate)
readonly hb2: Float32Array; // ffn hidden (up)
readonly logits: Float32Array;
readonly attScratch: Float32Array;
readonly qScratch: Float32Array; // one dequantised block
/** Counted so the simulator can show where the time goes. */
matmulCalls = 0;
matmulMacs = 0;
constructor(arena: Arena, dim: number, hidden: number, kvDim: number, vocab: number) {
this.x = arena.f32(dim, "ws.x");
this.xb = arena.f32(dim, "ws.xb");
this.xb2 = arena.f32(dim, "ws.xb2");
this.q = arena.f32(dim, "ws.q");
this.k = arena.f32(kvDim, "ws.k");
this.v = arena.f32(kvDim, "ws.v");
this.hb = arena.f32(hidden, "ws.hb");
this.hb2 = arena.f32(hidden, "ws.hb2");
this.logits = arena.f32(vocab, "ws.logits");
this.attScratch = arena.f32(Math.max(dim, 256), "ws.att");
this.qScratch = arena.f32(256, "ws.dequant"); // one Q4_K/Q6_K super-block
}
/**
* The single matmul entry point. Dispatches on how the weight is stored.
*
* Synchronous by design — see `tensor/webgpu.ts` for why the GPU path cannot
* live behind this signature, and why that is the honest answer rather than a
* limitation to route around.
*/
matmul(out: Float32Array, x: Float32Array, w: QuantWeight, n: number, d: number): void {
this.matmulCalls++;
this.matmulMacs += n * d;
if (w.kind === "f32") {
matmul(out, x, w.data, n, d);
return;
}
const bs = blockSize(w.qtype);
matmulQuant(out, x, w.raw, n, d, w.qtype, this.qScratch.subarray(0, bs));
}
resetCounters(): void {
this.matmulCalls = 0;
this.matmulMacs = 0;
}
}model/block.ts
Seven operations, twice through the residual stream. x is never copied.
/**
* mini-picoLM — step 10: the transformer block
* ============================================
*
* Seven operations, twice through the residual stream:
*
* x += attn(rmsnorm(x))
* x += ffn(rmsnorm(x))
*
* Everything before this file was a part; this is where they become a layer.
* Note that `x` is never copied — both sublayers read it, write into scratch,
* and add back. The residual stream is one buffer for the whole forward pass.
*/
import { rmsnorm } from "../tensor/rmsnorm.ts";
import { addInto } from "../tensor/matmul.ts";
import { ropeInPlace, type RopeTables } from "./rope.ts";
import { attention } from "./attention.ts";
import { feedForward } from "./mlp.ts";
import type { LayerWeights, ModelConfig } from "./llama.ts";
import type { Workspace } from "../runtime/workspace.ts";
import type { KVCache } from "../runtime/kv-cache.ts";
export function forwardBlock(
layer: number,
pos: number,
cfg: ModelConfig,
w: LayerWeights,
ws: Workspace,
cache: KVCache,
rope: RopeTables,
): void {
const { dim, nHeads, nKvHeads, headDim, kvDim, hidden } = cfg;
/* ---- attention ---- */
rmsnorm(ws.xb, ws.x, w.attnNorm);
ws.matmul(ws.q, ws.xb, w.attnQ, dim, dim);
ws.matmul(ws.k, ws.xb, w.attnK, dim, kvDim);
ws.matmul(ws.v, ws.xb, w.attnV, dim, kvDim);
// Rotate before caching: a cached key's position is fixed forever.
ropeInPlace(ws.q, nHeads, headDim, rope, pos);
ropeInPlace(ws.k, nKvHeads, headDim, rope, pos);
cache.write(layer, pos, ws.k, ws.v);
attention(ws.xb, ws.q, cache, layer, pos, { nHeads, nKvHeads, headDim }, ws.attScratch);
ws.matmul(ws.xb2, ws.xb, w.attnOut, dim, dim);
addInto(ws.x, ws.xb2);
/* ---- feed-forward ---- */
rmsnorm(ws.xb, ws.x, w.ffnNorm);
feedForward(ws.xb2, ws.xb, w.ffnGate, w.ffnUp, w.ffnDown, dim, hidden, ws);
addInto(ws.x, ws.xb2);
}model/llama.ts
The signature that is the whole course: forward(token, pos). One token per call is S01, pos indexing a persistent cache is S05, fp16-adjacent storage is S08, the online softmax is S12, nKvHeads < nHeads is S03.
/**
* mini-picoLM — the model
* =======================
*
* Config, weights, and `forward(token, pos)`.
*
* That signature is the whole course in one line. One token per call is the S01
* loop; `pos` indexing a persistent cache is S05; the weights arriving as
* quantised bytes is S08; the online softmax inside attention is S12;
* `nKvHeads < nHeads` is S03's GQA.
*
* There is no batch dimension anywhere, which is why there is no scheduler, no
* block manager and no continuous batching. One user, one tab.
*/
import { Arena } from "../memory/arena.ts";
import { rmsnorm } from "../tensor/rmsnorm.ts";
import { embedInto } from "./embedding.ts";
import { forwardBlock } from "./block.ts";
import { RopeTables } from "./rope.ts";
import { KVCache } from "../runtime/kv-cache.ts";
import { Workspace } from "../runtime/workspace.ts";
import { blockSize, rowSize, type QuantType } from "../format/quant.ts";
export type ModelConfig = {
dim: number;
hidden: number;
nHeads: number;
nKvHeads: number;
nLayers: number;
vocabSize: number;
maxSeq: number;
ropeFreqBase: number;
headDim: number;
kvDim: number;
};
export function makeConfig(
p: Omit<ModelConfig, "headDim" | "kvDim"> & Partial<Pick<ModelConfig, "headDim" | "kvDim">>,
): ModelConfig {
const headDim = p.headDim ?? p.dim / p.nHeads;
return { ...p, headDim, kvDim: p.kvDim ?? p.nKvHeads * headDim };
}
/** A weight is either a plain float view or quantised bytes plus its type. */
export type QuantWeight =
| { kind: "f32"; data: Float32Array }
| { kind: "quant"; raw: Uint8Array; qtype: QuantType };
export type LayerWeights = {
attnNorm: Float32Array;
ffnNorm: Float32Array;
attnQ: QuantWeight;
attnK: QuantWeight;
attnV: QuantWeight;
attnOut: QuantWeight;
ffnGate: QuantWeight;
ffnUp: QuantWeight;
ffnDown: QuantWeight;
};
export type ModelWeights = {
tokenEmbd: { raw: Uint8Array; qtype: QuantType };
outputNorm: Float32Array;
output: QuantWeight;
layers: LayerWeights[];
};
export class LlamaModel {
readonly cfg: ModelConfig;
readonly weights: ModelWeights;
readonly arena: Arena;
readonly cache: KVCache;
readonly ws: Workspace;
readonly rope: RopeTables;
constructor(cfg: ModelConfig, weights: ModelWeights, arena?: Arena) {
this.cfg = cfg;
this.weights = weights;
this.arena = arena ?? new Arena(LlamaModel.arenaBytes(cfg));
this.cache = new KVCache(this.arena, cfg.nLayers, cfg.maxSeq, cfg.kvDim);
this.ws = new Workspace(this.arena, cfg.dim, cfg.hidden, cfg.kvDim, cfg.vocabSize);
this.rope = new RopeTables(cfg.maxSeq, cfg.headDim, cfg.ropeFreqBase);
}
/** Exactly how much the arena needs — computed, not guessed. */
static arenaBytes(cfg: ModelConfig): number {
const kv = 2 * cfg.nLayers * cfg.maxSeq * cfg.kvDim * 4;
const ws = (4 * cfg.dim + 2 * cfg.kvDim + 2 * cfg.hidden + cfg.vocabSize + 512) * 4;
const slack = 64 * 1024; // alignment padding, one page of headroom
return kv + ws + slack;
}
get kvBytes(): number {
return this.cache.bytes;
}
reset(): void {
this.cache.reset();
this.ws.x.fill(0);
}
/** One token in, one row of logits out. */
forward(token: number, pos: number): Float32Array {
const { cfg, ws, weights } = this;
embedInto(
ws.x,
weights.tokenEmbd.raw,
token,
cfg.dim,
weights.tokenEmbd.qtype,
ws.qScratch.subarray(0, blockSize(weights.tokenEmbd.qtype)),
);
for (let l = 0; l < cfg.nLayers; l++) {
forwardBlock(l, pos, cfg, weights.layers[l], ws, this.cache, this.rope);
}
rmsnorm(ws.xb, ws.x, weights.outputNorm);
ws.matmul(ws.logits, ws.xb, weights.output, cfg.dim, cfg.vocabSize);
return ws.logits;
}
}
/** Bytes a quantised tensor of `n` weights occupies — used by the loader. */
export function weightBytes(qtype: QuantType, n: number): number {
return rowSize(qtype, n);
}model prints the config this page actually loaded, read back out of the GGUF header rather than out of a constant.
What is missing is the point
Continue
Steps 11–13, 17 — tokenizer, sampler, grammar, loop
tokenizer/bpe.ts
SentencePiece BPE: every token carries a score and the highest wins, unlike the ranked merges of S02's byte-level BPE. Spaces become ▁ and one is prepended; unknown characters fall back to <0xHH> byte tokens.
/**
* mini-picoLM — step 11: the BPE tokenizer
* ========================================
*
* SentencePiece-style BPE, the family Llama GGUF files carry. Every token has a
* **score**, and the merge with the highest score wins — not the lowest rank,
* which is the GPT/tiktoken convention.
*
* Three details that are easy to get wrong and hard to notice:
*
* * spaces become `▁` (U+2581) and one is *prepended*, so "Once" and " Once"
* tokenize identically;
* * anything not in the vocabulary falls back to a `<0xHH>` byte token, which
* is what makes the vocabulary total;
* * detokenization has to be **incremental and byte-aware**. A multi-byte
* character can span several tokens, so emitting each token's bytes as they
* arrive produces replacement characters in the UI. `StreamDecoder` holds
* partial sequences back.
*/
export const SPACE = "▁";
const BYTE_TOKEN = /^<0x([0-9A-Fa-f]{2})>$/;
export class Tokenizer {
readonly tokens: string[];
readonly scores: Float32Array;
readonly bosId: number;
readonly eosId: number;
private readonly lookup = new Map<string, number>();
private readonly byteOf = new Map<number, number>();
private readonly encoder = new TextEncoder();
constructor(tokens: string[], scores: Float32Array, bosId = 1, eosId = 2) {
this.tokens = tokens;
this.scores = scores;
this.bosId = bosId;
this.eosId = eosId;
for (let i = 0; i < tokens.length; i++) {
if (!this.lookup.has(tokens[i])) this.lookup.set(tokens[i], i);
const m = BYTE_TOKEN.exec(tokens[i]);
if (m) this.byteOf.set(i, parseInt(m[1], 16));
}
}
get vocabSize(): number {
return this.tokens.length;
}
/** One token per character, with a byte fallback for anything unknown. */
private seed(text: string): number[] {
const out: number[] = [];
for (const ch of text) {
const direct = this.lookup.get(ch);
if (direct !== undefined) {
out.push(direct);
continue;
}
for (const b of this.encoder.encode(ch)) {
const bt = this.lookup.get(`<0x${b.toString(16).toUpperCase().padStart(2, "0")}>`);
if (bt !== undefined) out.push(bt);
}
}
return out;
}
encode(text: string, addBos = true): number[] {
const norm = SPACE + text.replaceAll(" ", SPACE);
const parts = this.seed(norm);
// Greedy merge. Quadratic in the number of pieces, which is fine: it runs
// once per request on a short string.
for (;;) {
let bestScore = -Infinity;
let bestIdx = -1;
let bestId = -1;
for (let i = 0; i < parts.length - 1; i++) {
const merged = this.tokens[parts[i]] + this.tokens[parts[i + 1]];
const id = this.lookup.get(merged);
if (id !== undefined && this.scores[id] > bestScore) {
bestScore = this.scores[id];
bestIdx = i;
bestId = id;
}
}
if (bestIdx < 0) break;
parts.splice(bestIdx, 2, bestId);
}
return addBos ? [this.bosId, ...parts] : parts;
}
/** The raw bytes a token contributes. */
pieceBytes(id: number): Uint8Array {
const b = this.byteOf.get(id);
if (b !== undefined) return new Uint8Array([b]);
return this.encoder.encode(this.tokens[id].replaceAll(SPACE, " "));
}
/** Human-readable text for a token — what the grammar masks against. */
piece(id: number): string {
const b = this.byteOf.get(id);
if (b !== undefined) return String.fromCharCode(b);
return this.tokens[id].replaceAll(SPACE, " ");
}
decode(ids: number[], stripLeadingSpace = true): string {
let total = 0;
const parts = ids.map((i) => {
const p = this.pieceBytes(i);
total += p.length;
return p;
});
const buf = new Uint8Array(total);
let o = 0;
for (const p of parts) {
buf.set(p, o);
o += p.length;
}
const text = new TextDecoder().decode(buf);
return stripLeadingSpace && text.startsWith(" ") ? text.slice(1) : text;
}
}
/**
* Emits text only once the pending bytes form a complete character.
*
* `TextDecoder` with `{ stream: true }` does exactly this natively — it buffers
* an incomplete sequence internally and emits nothing until the continuation
* bytes arrive. Using it is both shorter and more correct than hand-rolling.
*/
export class StreamDecoder {
private readonly decoder = new TextDecoder("utf-8");
private readonly tok: Tokenizer;
constructor(tok: Tokenizer) {
this.tok = tok;
}
push(id: number): string {
return this.decoder.decode(this.tok.pieceBytes(id), { stream: true });
}
/** Flush anything still buffered at end of stream. */
finish(): string {
return this.decoder.decode(new Uint8Array(0));
}
}Edit the text and run it. The command asserts its own round trip, so a tokenizer that loses information says so.
runtime/sampler.ts
Greedy at temperature 0, otherwise temperature → softmax → top-k → top-p. The >= in cutoff is S04's +1: the token that crosses the threshold is kept; without it, a distribution whose top token already exceeds top_p would leave nothing selected and divide by zero.
/**
* mini-picoLM — step 12: sampling
* ===============================
*
* Greedy at temperature 0, otherwise temperature → softmax → top-k → top-p.
*
* Two things carried over from S04 that matter in production and are usually
* got wrong:
*
* **The `+1`.** Top-p keeps the *smallest set whose cumulative mass reaches
* p*, which means the token that crosses the threshold is included. Drop it
* and a distribution whose top token already has p=0.95 keeps nothing at
* `top_p=0.9`, and the renormalisation divides by zero.
*
* **The RNG is per-sampler, not global.** With a shared generator a request's
* output depends on what else was sampled first, which makes bug reports
* impossible to reproduce. This one is seeded and owns its state.
*/
/** mulberry32 — small, fast, and reproducible across engines. */
export function makeRng(seed: number): () => number {
let a = seed >>> 0;
return function next(): number {
a = (a + 0x6d2b79f5) >>> 0;
let t = a;
t = Math.imul(t ^ (t >>> 15), t | 1);
t ^= t + Math.imul(t ^ (t >>> 7), t | 61);
return ((t ^ (t >>> 14)) >>> 0) / 4294967296;
};
}
export type SamplerOptions = {
temperature?: number;
topK?: number;
topP?: number;
seed?: number;
};
export class Sampler {
temperature: number;
topK: number;
topP: number;
private rng: () => number;
private order: Int32Array | null = null;
private probs: Float32Array | null = null;
constructor(opts: SamplerOptions = {}) {
this.temperature = opts.temperature ?? 0.8;
this.topK = opts.topK ?? 0;
this.topP = opts.topP ?? 0.9;
this.rng = makeRng(opts.seed ?? 0);
}
reseed(seed: number): void {
this.rng = makeRng(seed);
}
/** Number of tokens that survive the filters — what the simulator plots. */
reachable(logits: Float32Array): number {
if (this.temperature <= 0) return 1;
const { probs, order } = this.prepare(logits);
return this.cutoff(probs, order);
}
sample(logits: Float32Array): number {
if (this.temperature <= 0) {
let best = 0;
for (let i = 1; i < logits.length; i++) if (logits[i] > logits[best]) best = i;
return best;
}
const { probs, order } = this.prepare(logits);
const keep = this.cutoff(probs, order);
let mass = 0;
for (let i = 0; i < keep; i++) mass += probs[order[i]];
let r = this.rng() * mass;
for (let i = 0; i < keep; i++) {
r -= probs[order[i]];
if (r <= 0) return order[i];
}
return order[keep - 1];
}
/** Softmax at temperature, plus indices sorted by descending probability. */
private prepare(logits: Float32Array): { probs: Float32Array; order: Int32Array } {
const n = logits.length;
if (!this.probs || this.probs.length !== n) {
this.probs = new Float32Array(n);
this.order = new Int32Array(n);
}
const probs = this.probs;
const order = this.order!;
let max = -Infinity;
for (let i = 0; i < n; i++) {
const v = logits[i] / this.temperature;
probs[i] = v;
if (v > max) max = v;
}
let sum = 0;
for (let i = 0; i < n; i++) {
const e = Math.exp(probs[i] - max);
probs[i] = e;
sum += e;
}
const inv = 1 / sum;
for (let i = 0; i < n; i++) {
probs[i] *= inv;
order[i] = i;
}
// A full sort is O(V log V) per token — 130k entries in a real vocabulary.
// Production kernels use a radix-select on the threshold instead; this is
// the honest simple version, and the cost is visible in the simulator.
const idx = Array.from(order);
idx.sort((a, b) => probs[b] - probs[a]);
order.set(idx);
return { probs, order };
}
/** How many of the sorted tokens survive top-k and top-p. */
private cutoff(probs: Float32Array, order: Int32Array): number {
const n = order.length;
const limit = this.topK > 0 ? Math.min(this.topK, n) : n;
if (this.topP >= 1) return limit;
let cum = 0;
for (let i = 0; i < limit; i++) {
cum += probs[order[i]];
// >= keeps the token that crosses the threshold. This is the `+1`.
if (cum >= this.topP) return i + 1;
}
return limit;
}
}runtime/grammar.ts
A pushdown automaton over JSON, lifted to the token level. The container stack is what makes it pushdown rather than finite-state. Nesting is unbounded, so no fixed set of states can remember whether the next } is legal.
/**
* mini-picoLM — step 17: grammar constraints
* ==========================================
*
* A pushdown automaton over JSON, lifted to the token level: a token is legal
* only if *every* character it would emit keeps the machine alive.
*
* The container stack is what makes this pushdown rather than finite-state.
* JSON nesting is unbounded, so no fixed set of states can remember whether the
* next `}` is legal — a stack can.
*
* Illegal tokens are set to -Infinity before the sampler runs, so invalid
* output is not discouraged, it is unreachable at any temperature.
*
* What this does *not* guarantee is completion. A token budget can cut the
* object off mid-structure, and a truncated `{"name": "ad` is as unusable to
* the caller as invalid JSON. Check `accepting()` before you return.
*/
const WS = " \t\n\r";
const DIGITS = "0123456789";
const LITERALS: Record<string, string> = { t: "rue", f: "alse", n: "ull" };
export type GrammarState = {
node: string;
stack: readonly string[];
/** Characters a literal still owes, or "d" once a number has seen a digit. */
buf: string;
};
export class JsonGrammar {
private readonly maxDepth: number;
constructor(maxDepth = 16) {
this.maxDepth = maxDepth;
}
start(): GrammarState {
return { node: "value", stack: [], buf: "" };
}
step(s: GrammarState, ch: string): GrammarState | null {
const { node, stack, buf } = s;
switch (node) {
case "value":
if (WS.includes(ch)) return s;
if (ch === "{")
return stack.length < this.maxDepth
? { node: "objOpen", stack: [...stack, "o"], buf: "" }
: null;
if (ch === "[")
return stack.length < this.maxDepth
? { node: "arrOpen", stack: [...stack, "a"], buf: "" }
: null;
if (ch === '"') return { node: "string", stack, buf: "" };
if (ch === "-") return { node: "number", stack, buf: "" };
if (DIGITS.includes(ch)) return { node: "number", stack, buf: "d" };
if (ch in LITERALS) return { node: "literal", stack, buf: LITERALS[ch] };
return null;
case "literal":
if (!buf || ch !== buf[0]) return null;
return buf.length === 1
? { node: "after", stack, buf: "" }
: { node: "literal", stack, buf: buf.slice(1) };
case "number":
if (DIGITS.includes(ch)) return { node: "number", stack, buf: "d" };
if (".eE+-".includes(ch) && buf === "d") return { node: "number", stack, buf: "d" };
return buf === "d" ? this.step({ node: "after", stack, buf: "" }, ch) : null;
case "string":
if (ch === "\\") return { node: "stringEsc", stack, buf: "" };
if (ch === '"') return { node: "after", stack, buf: "" };
return "\n\r".includes(ch) ? null : s;
case "stringEsc":
return '"\\/bfnrtu'.includes(ch) ? { node: "string", stack, buf: "" } : null;
case "objOpen":
if (WS.includes(ch)) return s;
if (ch === '"') return { node: "key", stack, buf: "" };
return ch === "}" ? this.close(stack, "o") : null;
case "key":
if (ch === '"') return { node: "keyEnd", stack, buf: "" };
return "\n\r".includes(ch) ? null : s;
case "keyEnd":
if (WS.includes(ch)) return s;
return ch === ":" ? { node: "value", stack, buf: "" } : null;
case "commaObj":
if (WS.includes(ch)) return s;
return ch === '"' ? { node: "key", stack, buf: "" } : null;
case "arrOpen":
if (WS.includes(ch)) return s;
if (ch === "]") return this.close(stack, "a");
return this.step({ node: "value", stack, buf: "" }, ch);
case "after": {
if (WS.includes(ch)) return s;
if (stack.length === 0) return null;
const top = stack[stack.length - 1];
if (ch === ",")
return top === "o"
? { node: "commaObj", stack, buf: "" }
: { node: "value", stack, buf: "" };
if (ch === "}" && top === "o") return this.close(stack, "o");
if (ch === "]" && top === "a") return this.close(stack, "a");
return null;
}
default:
return null;
}
}
private close(stack: readonly string[], kind: string): GrammarState | null {
if (!stack.length || stack[stack.length - 1] !== kind) return null;
return { node: "after", stack: stack.slice(0, -1), buf: "" };
}
accepting(s: GrammarState | null): boolean {
if (!s || s.stack.length) return false;
return s.node === "after" || (s.node === "number" && s.buf === "d");
}
/** Feed a whole token piece. Null if any character leaves the grammar. */
advance(s: GrammarState | null, text: string): GrammarState | null {
let cur = s;
for (const ch of text) {
if (!cur) return null;
cur = this.step(cur, ch);
}
return cur;
}
/**
* Set every grammar-illegal logit to -Infinity, in place.
*
* Cost is O(vocab × token length) per step, which is why production engines
* compile and cache these masks per state. Returns how many survived.
*/
maskLogits(logits: Float32Array, s: GrammarState, pieces: string[]): number {
let legal = 0;
for (let i = 0; i < pieces.length; i++) {
if (this.advance(s, pieces[i]) === null) logits[i] = -Infinity;
else legal++;
}
return legal;
}
}The grammar group of the self-test checks the automaton against JSON.parse on nine strings, including the ones designed to be almost valid.
runtime/generation.ts
The async generator. It yields after every token so the event loop can paint, and the caller stops it by not asking for the next value.
/**
* mini-picoLM — step 13: generation
* =================================
*
* Prefill, then the loop. After every other step in this list, it is still the
* while-loop S01 opened with.
*
* The one thing that is genuinely different in a browser is that **you may not
* block the main thread**. A synchronous loop over 200 tokens freezes the tab:
* no repaint, no scrolling, no cancel button. So `generate` is an async
* generator that yields after each token, which gives the event loop a chance
* to paint and lets the caller stop simply by not asking for the next value.
*
* That is the browser's version of the cancellation story in S19 — and it is
* why streaming is not a UI nicety here but a structural requirement.
*/
import type { LlamaModel } from "../model/llama.ts";
import type { Tokenizer } from "../tokenizer/bpe.ts";
import { StreamDecoder } from "../tokenizer/bpe.ts";
import type { Sampler } from "./sampler.ts";
import { JsonGrammar, type GrammarState } from "./grammar.ts";
export type GenerateOptions = {
prompt: string;
maxTokens?: number;
json?: boolean;
/** Yield to the event loop every N tokens. 1 keeps the UI perfectly live. */
yieldEvery?: number;
signal?: { aborted: boolean };
};
export type TokenEvent = {
token: number;
text: string;
index: number;
position: number;
/** ms spent inside forward() for this token */
ms: number;
legalTokens?: number;
};
export type GenerateResult = {
text: string;
promptTokens: number;
outputTokens: number;
prefillMs: number;
decodeMs: number;
tokensPerSecond: number;
stopReason: "eos" | "maxTokens" | "contextFull" | "aborted" | "grammarComplete";
grammarValid?: boolean;
};
const now = (): number =>
typeof performance !== "undefined" ? performance.now() : Date.now();
/** Hand control back so the browser can paint. A no-op cost in Node. */
const breathe = (): Promise<void> =>
new Promise((resolve) => {
setTimeout(resolve, 0);
});
export async function* generate(
model: LlamaModel,
tok: Tokenizer,
sampler: Sampler,
opts: GenerateOptions,
): AsyncGenerator<TokenEvent, GenerateResult, void> {
const maxTokens = opts.maxTokens ?? 64;
const yieldEvery = opts.yieldEvery ?? 1;
const ids = tok.encode(opts.prompt, true);
const dec = new StreamDecoder(tok);
const grammar = opts.json ? new JsonGrammar() : null;
let gstate: GrammarState | null = grammar ? grammar.start() : null;
const pieces = grammar ? tok.tokens.map((_, i) => tok.piece(i)) : [];
/* ---- prefill: one position at a time, exactly like decode ---- */
const t0 = now();
let logits = model.ws.logits;
for (let i = 0; i < ids.length; i++) {
logits = model.forward(ids[i], i);
}
const prefillMs = now() - t0;
let pos = ids.length;
let produced = 0;
let out = "";
let stopReason: GenerateResult["stopReason"] = "maxTokens";
const decodeStart = now();
while (produced < maxTokens) {
if (opts.signal?.aborted) {
stopReason = "aborted";
break;
}
if (pos >= model.cfg.maxSeq) {
stopReason = "contextFull";
break;
}
let legal: number | undefined;
if (grammar && gstate) {
legal = grammar.maskLogits(logits, gstate, pieces);
if (legal === 0) {
stopReason = "grammarComplete";
break;
}
}
const next = sampler.sample(logits);
if (next === tok.eosId) {
stopReason = "eos";
break;
}
if (grammar && gstate) {
const advanced = grammar.advance(gstate, pieces[next]);
if (advanced) gstate = advanced;
}
const text = dec.push(next);
out += text;
const tf = now();
logits = model.forward(next, pos);
const ms = now() - tf;
yield { token: next, text, index: produced, position: pos, ms, legalTokens: legal };
pos++;
produced++;
if (grammar && grammar.accepting(gstate)) {
stopReason = "grammarComplete";
break;
}
if (produced % yieldEvery === 0) await breathe();
}
out += dec.finish();
const decodeMs = now() - decodeStart;
return {
text: out,
promptTokens: ids.length,
outputTokens: produced,
prefillMs,
decodeMs,
tokensPerSecond: decodeMs > 0 ? (produced / decodeMs) * 1000 : 0,
stopReason,
grammarValid: grammar ? grammar.accepting(gstate) : undefined,
};
}
/** Convenience wrapper for callers that just want the string. */
export async function generateText(
model: LlamaModel,
tok: Tokenizer,
sampler: Sampler,
opts: GenerateOptions,
): Promise<GenerateResult> {
const it = generate(model, tok, sampler, opts);
for (;;) {
const r = await it.next();
if (r.done) return r.value;
}
}Streaming, here, now. Add --json to constrain the output and the command will run JSON.parse over whatever came back.
Masking guarantees valid, not complete
{"name": "ad is as unusable to the caller as invalid JSON. That is why generate reports grammarValid separately from stopReason.Step 8
Step 14 — GGUF, and the entry point
In C, loading a model is mmap. In a browser there is no mmap, but an ArrayBuffer from a file input or a fetch plays the same role: one region, and a tensor is a subarray at a known offset. The subarray calls are the equivalent of picoLM's pointer arithmetic, and they are equally free.
format/gguf.ts
Reader and writer. The writer exists so the page can synthesise a valid model and exercise the reader with no download. The loader is therefore tested on every run rather than only when someone has a 640 MB file to hand.
/**
* mini-picoLM — step 14: the GGUF reader
* ======================================
*
* Parse the container, then hand out `Uint8Array` views onto the tensor data.
* Nothing is copied and nothing is dequantised at load time.
*
* In C this is `mmap`. In a browser there is no mmap — but an `ArrayBuffer`
* from `File.arrayBuffer()` or a streamed `fetch` plays the same role: one
* region, and a tensor is a subarray at a known offset. The `subarray` calls
* below are the equivalent of picoLM's pointer arithmetic, and they are equally
* free.
*
* The writer exists so the S22 simulator can synthesise a valid model in the
* page and exercise this reader with no download.
*/
import { GGUF_TYPE, rowSize, type QuantType } from "./quant.ts";
export const GGUF_MAGIC = 0x46554747; // "GGUF"
export const DEFAULT_ALIGNMENT = 32;
const U8 = 0,
I8 = 1,
U16 = 2,
I16 = 3,
U32 = 4,
I32 = 5,
F32 = 6,
BOOL = 7,
STRING = 8,
ARRAY = 9,
U64 = 10,
I64 = 11,
F64 = 12;
export type MetaValue = number | boolean | string | MetaValue[];
export type TensorInfo = {
name: string;
dims: number[];
qtype: QuantType;
offset: number;
nElements: number;
nBytes: number;
};
class Reader {
pos = 0;
readonly dv: DataView;
readonly bytes: Uint8Array;
constructor(dv: DataView, bytes: Uint8Array) {
this.dv = dv;
this.bytes = bytes;
}
u32(): number {
const v = this.dv.getUint32(this.pos, true);
this.pos += 4;
return v;
}
i32(): number {
const v = this.dv.getInt32(this.pos, true);
this.pos += 4;
return v;
}
/** GGUF counts are u64; JS numbers hold them exactly below 2^53. */
u64(): number {
const lo = this.dv.getUint32(this.pos, true);
const hi = this.dv.getUint32(this.pos + 4, true);
this.pos += 8;
if (hi > 0x1fffff) throw new RangeError("GGUF value exceeds Number.MAX_SAFE_INTEGER");
return hi * 2 ** 32 + lo;
}
str(): string {
const n = this.u64();
const s = new TextDecoder().decode(this.bytes.subarray(this.pos, this.pos + n));
this.pos += n;
return s;
}
value(t: number): MetaValue {
switch (t) {
case U8: return this.bytes[this.pos++];
case I8: return this.dv.getInt8(this.pos++);
case U16: { const v = this.dv.getUint16(this.pos, true); this.pos += 2; return v; }
case I16: { const v = this.dv.getInt16(this.pos, true); this.pos += 2; return v; }
case U32: return this.u32();
case I32: return this.i32();
case F32: { const v = this.dv.getFloat32(this.pos, true); this.pos += 4; return v; }
case F64: { const v = this.dv.getFloat64(this.pos, true); this.pos += 8; return v; }
case BOOL: return this.bytes[this.pos++] !== 0;
case U64: return this.u64();
case I64: return this.u64();
case STRING: return this.str();
case ARRAY: {
const at = this.u32();
const n = this.u64();
const out: MetaValue[] = new Array(n);
for (let i = 0; i < n; i++) out[i] = this.value(at);
return out;
}
default:
throw new Error(`unknown GGUF metadata type ${t}`);
}
}
}
export class GGUFFile {
readonly metadata = new Map<string, MetaValue>();
readonly tensors = new Map<string, TensorInfo>();
readonly dataStart: number;
private readonly bytes: Uint8Array;
constructor(buffer: ArrayBuffer) {
this.bytes = new Uint8Array(buffer);
const r = new Reader(new DataView(buffer), this.bytes);
if (r.u32() !== GGUF_MAGIC) {
throw new Error("not a GGUF file (bad magic) — is this really a .gguf?");
}
const version = r.u32();
if (version < 2) throw new Error(`GGUF version ${version} is too old; need >= 2`);
const nTensors = r.u64();
const nKv = r.u64();
for (let i = 0; i < nKv; i++) {
const key = r.str();
this.metadata.set(key, r.value(r.u32()));
}
const infos: TensorInfo[] = [];
for (let i = 0; i < nTensors; i++) {
const name = r.str();
const nDims = r.u32();
const dims: number[] = [];
for (let d = 0; d < nDims; d++) dims.push(r.u64());
const raw = r.u32();
const qtype = GGUF_TYPE[raw];
const offset = r.u64();
if (!qtype) {
throw new Error(
`tensor "${name}" uses GGUF type ${raw}, which mini-picoLM does not read. ` +
`Supported: ${Object.values(GGUF_TYPE).join(", ")}.`,
);
}
const nElements = dims.reduce((a, b) => a * b, 1);
infos.push({ name, dims, qtype, offset, nElements, nBytes: rowSize(qtype, nElements) });
}
const align = Number(this.metadata.get("general.alignment") ?? DEFAULT_ALIGNMENT);
this.dataStart = Math.ceil(r.pos / align) * align;
for (const t of infos) {
if (this.dataStart + t.offset + t.nBytes > this.bytes.length) {
throw new Error(`tensor "${t.name}" runs past the end of the file`);
}
this.tensors.set(t.name, t);
}
}
has(name: string): boolean {
return this.tensors.has(name);
}
info(name: string): TensorInfo {
const t = this.tensors.get(name);
if (!t) throw new Error(`no tensor named "${name}" in this GGUF file`);
return t;
}
/** The tensor's bytes, as a view. No copy — this is the mmap equivalent. */
raw(name: string): Uint8Array {
const t = this.info(name);
return this.bytes.subarray(this.dataStart + t.offset, this.dataStart + t.offset + t.nBytes);
}
num(key: string, fallback?: number): number {
const v = this.metadata.get(key);
if (typeof v === "number") return v;
if (fallback !== undefined) return fallback;
throw new Error(`GGUF metadata is missing "${key}"`);
}
str(key: string, fallback?: string): string {
const v = this.metadata.get(key);
if (typeof v === "string") return v;
if (fallback !== undefined) return fallback;
throw new Error(`GGUF metadata is missing "${key}"`);
}
get totalTensorBytes(): number {
let n = 0;
for (const t of this.tensors.values()) n += t.nBytes;
return n;
}
}
/* ------------------------------------------------------------- writing */
type WriteTensor = { data: Float32Array; dims: number[] };
function pushString(out: number[], s: string): void {
const b = new TextEncoder().encode(s);
pushU64(out, b.length);
for (const x of b) out.push(x);
}
function pushU32(out: number[], v: number): void {
out.push(v & 0xff, (v >>> 8) & 0xff, (v >>> 16) & 0xff, (v >>> 24) & 0xff);
}
function pushU64(out: number[], v: number): void {
const lo = v >>> 0;
const hi = Math.floor(v / 2 ** 32);
pushU32(out, lo);
pushU32(out, hi);
}
function pushValue(out: number[], v: MetaValue, tagged = true): void {
if (typeof v === "boolean") {
if (tagged) pushU32(out, BOOL);
out.push(v ? 1 : 0);
} else if (typeof v === "number") {
if (Number.isInteger(v) && v >= 0 && v < 2 ** 31) {
if (tagged) pushU32(out, U32);
pushU32(out, v);
} else {
if (tagged) pushU32(out, F32);
const b = new Uint8Array(4);
new DataView(b.buffer).setFloat32(0, v, true);
out.push(...b);
}
} else if (typeof v === "string") {
if (tagged) pushU32(out, STRING);
pushString(out, v);
} else {
if (tagged) pushU32(out, ARRAY);
const inner = v.length ? v[0] : "";
const at =
typeof inner === "string"
? STRING
: typeof inner === "boolean"
? BOOL
: Number.isInteger(inner as number) && (inner as number) >= 0
? U32
: F32;
pushU32(out, at);
pushU64(out, v.length);
for (const item of v) pushValue(out, item, false);
}
}
/**
* Write a valid GGUF file into a fresh ArrayBuffer.
*
* Only F32 tensors are written — the point is to exercise the container and the
* loader, not to reimplement a quantiser. The quantised paths are checked
* against hand-built blocks in `selftest.ts`.
*/
export function writeGGUF(
metadata: Record<string, MetaValue>,
tensors: Record<string, WriteTensor>,
): ArrayBuffer {
const align = Number(metadata["general.alignment"] ?? DEFAULT_ALIGNMENT);
const names = Object.keys(tensors);
const head: number[] = [];
pushU32(head, GGUF_MAGIC);
pushU32(head, 3);
pushU64(head, names.length);
pushU64(head, Object.keys(metadata).length);
for (const [k, v] of Object.entries(metadata)) {
pushString(head, k);
pushValue(head, v);
}
let offset = 0;
const blobs: { bytes: Uint8Array; padded: number }[] = [];
for (const name of names) {
const { data, dims } = tensors[name];
pushString(head, name);
pushU32(head, dims.length);
// GGUF dims are fastest-varying first — the reverse of a row-major shape.
for (const d of [...dims].reverse()) pushU64(head, d);
pushU32(head, 0); // F32
pushU64(head, offset);
const bytes = new Uint8Array(data.buffer, data.byteOffset, data.byteLength).slice();
const padded = Math.ceil(bytes.length / align) * align;
blobs.push({ bytes, padded });
offset += padded;
}
const headerLen = Math.ceil(head.length / align) * align;
const total = headerLen + offset;
const buf = new ArrayBuffer(total);
const view = new Uint8Array(buf);
view.set(head, 0);
let o = headerLen;
for (const b of blobs) {
view.set(b.bytes, o);
o += b.padded;
}
return buf;
}main.ts
loadFromGGUF for a real model, createToyModel for one built in memory. The toy path goes through the file format rather than around it, so it exercises the reader too.
/**
* mini-picoLM — the public entry point
* ====================================
*
* Two ways in:
*
* `loadFromGGUF(buffer)` — a real model. Feed it an `ArrayBuffer` from a file
* input, a `fetch`, or drag-and-drop. Nothing is copied; the tensors stay
* quantised and are read in place.
*
* `createToyModel()` — a small, randomly-weighted model built in memory,
* through a genuine GGUF file this module writes itself. The output is noise,
* but every code path is the real one, which is what makes the whole engine
* testable in a browser with no download.
*/
import { Arena } from "./memory/arena.ts";
import { GGUFFile, writeGGUF } from "./format/gguf.ts";
import { dequantizeRow, rowSize } from "./format/quant.ts";
import {
LlamaModel,
makeConfig,
type LayerWeights,
type ModelConfig,
type ModelWeights,
type QuantWeight,
} from "./model/llama.ts";
import { Tokenizer } from "./tokenizer/bpe.ts";
import { Sampler, type SamplerOptions } from "./runtime/sampler.ts";
export type Engine = {
model: LlamaModel;
tokenizer: Tokenizer;
sampler: Sampler;
config: ModelConfig;
/** Bytes of weights, still quantised, referenced not copied. */
weightBytes: number;
/** Bytes the arena actually reserved (KV cache + workspace). */
arenaBytes: number;
source: string;
};
function asWeight(f: GGUFFile, name: string): QuantWeight {
const info = f.info(name);
if (info.qtype !== "F32") return { kind: "quant", raw: f.raw(name), qtype: info.qtype };
const copy = f.raw(name).slice(); // copy: the view is not guaranteed 4-aligned
return { kind: "f32", data: new Float32Array(copy.buffer, 0, info.nElements) };
}
export type LoadOptions = {
maxSeq?: number;
sampler?: SamplerOptions;
};
export function loadFromGGUF(buffer: ArrayBuffer, opts: LoadOptions = {}): Engine {
const f = new GGUFFile(buffer);
const arch = f.str("general.architecture", "llama");
const p = `${arch}.`;
const dim = f.num(`${p}embedding_length`);
const nHeads = f.num(`${p}attention.head_count`);
const trained = f.num(`${p}context_length`, 2048);
const tokensMeta = f.metadata.get("tokenizer.ggml.tokens");
if (!Array.isArray(tokensMeta)) {
throw new Error("this GGUF file carries no tokenizer.ggml.tokens array");
}
const tokens = tokensMeta as string[];
const cfg = makeConfig({
dim,
hidden: f.num(`${p}feed_forward_length`),
nHeads,
nKvHeads: f.num(`${p}attention.head_count_kv`, nHeads),
nLayers: f.num(`${p}block_count`),
vocabSize: f.num(`${p}vocab_size`, tokens.length),
maxSeq: Math.min(opts.maxSeq ?? trained, trained),
ropeFreqBase: f.num(`${p}rope.freq_base`, 10000),
});
const layers: LayerWeights[] = [];
for (let l = 0; l < cfg.nLayers; l++) {
layers.push({
attnNorm: dequantToFloats(f, `blk.${l}.attn_norm.weight`),
ffnNorm: dequantToFloats(f, `blk.${l}.ffn_norm.weight`),
attnQ: asWeight(f, `blk.${l}.attn_q.weight`),
attnK: asWeight(f, `blk.${l}.attn_k.weight`),
attnV: asWeight(f, `blk.${l}.attn_v.weight`),
attnOut: asWeight(f, `blk.${l}.attn_output.weight`),
ffnGate: asWeight(f, `blk.${l}.ffn_gate.weight`),
ffnUp: asWeight(f, `blk.${l}.ffn_up.weight`),
ffnDown: asWeight(f, `blk.${l}.ffn_down.weight`),
});
}
const embdInfo = f.info("token_embd.weight");
const weights: ModelWeights = {
tokenEmbd: { raw: f.raw("token_embd.weight"), qtype: embdInfo.qtype },
outputNorm: dequantToFloats(f, "output_norm.weight"),
// Many models tie the output projection to the embedding table.
output: asWeight(f, f.has("output.weight") ? "output.weight" : "token_embd.weight"),
layers,
};
const scoresMeta = f.metadata.get("tokenizer.ggml.scores");
const scores = new Float32Array(tokens.length);
if (Array.isArray(scoresMeta)) {
for (let i = 0; i < tokens.length; i++) scores[i] = Number(scoresMeta[i] ?? 0);
}
const tokenizer = new Tokenizer(
tokens,
scores,
f.num("tokenizer.ggml.bos_token_id", 1),
f.num("tokenizer.ggml.eos_token_id", 2),
);
const arena = new Arena(LlamaModel.arenaBytes(cfg));
const model = new LlamaModel(cfg, weights, arena);
return {
model,
tokenizer,
sampler: new Sampler(opts.sampler),
config: cfg,
weightBytes: f.totalTensorBytes,
arenaBytes: arena.capacity,
source: `${arch} · ${f.tensors.size} tensors`,
};
}
/**
* Dequantise a whole (small) tensor to floats. Used only for norm weights,
* which are read every layer and are a few kilobytes in total.
*/
function dequantToFloats(f: GGUFFile, name: string): Float32Array {
const info = f.info(name);
const raw = f.raw(name);
if (info.qtype === "F32") {
// `raw` is a view at an arbitrary byte offset, and Float32Array can only
// alias a 4-aligned one — so copy these few hundred bytes rather than
// depend on the file's padding.
const copy = raw.slice();
return new Float32Array(copy.buffer, 0, info.nElements);
}
return dequantizeRow(raw, info.nElements, info.qtype);
}
/* --------------------------------------------------- the in-memory model */
export type ToyOptions = {
nLayers?: number;
dim?: number;
nHeads?: number;
nKvHeads?: number;
hidden?: number;
maxSeq?: number;
seed?: number;
};
/** A deterministic PRNG so the toy model is identical everywhere. */
function toyRng(seed: number): () => number {
let a = seed >>> 0;
return () => {
a = (a + 0x6d2b79f5) >>> 0;
let t = a;
t = Math.imul(t ^ (t >>> 15), t | 1);
t ^= t + Math.imul(t ^ (t >>> 7), t | 61);
return (((t ^ (t >>> 14)) >>> 0) / 4294967296 - 0.5) * 2;
};
}
export const TOY_VOCAB: string[] = [
"<unk>",
"<s>",
"</s>",
...Array.from({ length: 256 }, (_, b) => `<0x${b.toString(16).toUpperCase().padStart(2, "0")}>`),
"▁",
"▁the",
"▁cat",
"▁sat",
"▁on",
"▁mat",
"▁and",
"a",
"b",
"c",
"{",
"}",
'"',
":",
",",
"[",
"]",
"▁1",
"▁2",
"true",
"false",
"▁name",
"▁age",
"▁ok",
];
/**
* Build a small model, serialise it as a real GGUF file, and load it back
* through the real reader.
*
* Going through the file format rather than constructing the weights directly
* is the point: it means the browser demo exercises `format/gguf.ts` on every
* run, so the loader cannot rot.
*/
export function createToyModel(
opts: ToyOptions = {},
): Engine & { ggufBytes: number; ggufBuffer: ArrayBuffer } {
const nLayers = opts.nLayers ?? 2;
const dim = opts.dim ?? 64;
const nHeads = opts.nHeads ?? 4;
const nKvHeads = opts.nKvHeads ?? 2;
const hidden = opts.hidden ?? 128;
const maxSeq = opts.maxSeq ?? 128;
const vocab = TOY_VOCAB.length;
const headDim = dim / nHeads;
const kvDim = nKvHeads * headDim;
const rnd = toyRng(opts.seed ?? 1);
const mk = (n: number, scale = 0.08) => {
const a = new Float32Array(n);
for (let i = 0; i < n; i++) a[i] = rnd() * scale;
return a;
};
const ones = (n: number) => new Float32Array(n).fill(1);
const tensors: Record<string, { data: Float32Array; dims: number[] }> = {
"token_embd.weight": { data: mk(vocab * dim), dims: [vocab, dim] },
"output_norm.weight": { data: ones(dim), dims: [dim] },
"output.weight": { data: mk(vocab * dim), dims: [vocab, dim] },
};
for (let l = 0; l < nLayers; l++) {
tensors[`blk.${l}.attn_norm.weight`] = { data: ones(dim), dims: [dim] };
tensors[`blk.${l}.ffn_norm.weight`] = { data: ones(dim), dims: [dim] };
tensors[`blk.${l}.attn_q.weight`] = { data: mk(dim * dim), dims: [dim, dim] };
tensors[`blk.${l}.attn_k.weight`] = { data: mk(kvDim * dim), dims: [kvDim, dim] };
tensors[`blk.${l}.attn_v.weight`] = { data: mk(kvDim * dim), dims: [kvDim, dim] };
tensors[`blk.${l}.attn_output.weight`] = { data: mk(dim * dim), dims: [dim, dim] };
tensors[`blk.${l}.ffn_gate.weight`] = { data: mk(hidden * dim), dims: [hidden, dim] };
tensors[`blk.${l}.ffn_up.weight`] = { data: mk(hidden * dim), dims: [hidden, dim] };
tensors[`blk.${l}.ffn_down.weight`] = { data: mk(dim * hidden), dims: [dim, hidden] };
}
const buffer = writeGGUF(
{
"general.architecture": "llama",
"general.alignment": 32,
"llama.embedding_length": dim,
"llama.feed_forward_length": hidden,
"llama.attention.head_count": nHeads,
"llama.attention.head_count_kv": nKvHeads,
"llama.block_count": nLayers,
"llama.context_length": maxSeq,
"llama.vocab_size": vocab,
"llama.rope.freq_base": 10000,
"tokenizer.ggml.tokens": TOY_VOCAB,
"tokenizer.ggml.scores": TOY_VOCAB.map((_, i) => -i),
"tokenizer.ggml.bos_token_id": 1,
"tokenizer.ggml.eos_token_id": 2,
},
tensors,
);
const engine = loadFromGGUF(buffer, { maxSeq, sampler: { temperature: 0.9, topP: 0.9 } });
return {
...engine,
source: "toy model (in-memory GGUF)",
ggufBytes: buffer.byteLength,
ggufBuffer: buffer,
};
}
export { rowSize };
export * from "./runtime/generation.ts";
export { Sampler } from "./runtime/sampler.ts";
export { Tokenizer, StreamDecoder } from "./tokenizer/bpe.ts";
export { JsonGrammar } from "./runtime/grammar.ts";
export { Arena } from "./memory/arena.ts";
export { GGUFFile, writeGGUF } from "./format/gguf.ts";
export { LlamaModel } from "./model/llama.ts";
export { benchmark, probeWebGPU, hasWebGPU } from "./tensor/webgpu.ts";gguf parses the file this page wrote at startup and lists what it found: magic, alignment, and the tensor index.
Step 9
Step 18 — WebGPU
tensor/webgpu.ts
A real compute shader: one workgroup per output row, 64 lanes striding through it, then a tree reduction. Everything is feature-detected, because navigator.gpu is absent during SSR, absent in several shipping browsers, and gated behind a flag in others.
/**
* mini-picoLM — step 18: WebGPU acceleration
* ==========================================
*
* A real compute shader for the operation that dominates decode: matrix-vector
* multiply. One workgroup per output row, a 64-lane tree reduction inside it.
*
* **Why this is not inside the token loop.** Reading a result back from the GPU
* is asynchronous — `mapAsync` returns a promise — while `Workspace.matmul` is
* synchronous and called seven times per layer. Making the forward pass async
* would mean awaiting dozens of GPU round trips per token, and a round trip
* costs more than the matmul saves at these sizes. So the CPU path stays the
* reference, and the GPU path is exposed as an explicit async call used for
* prefill-shaped work and for the benchmark the S22 simulator runs.
*
* That is not a cop-out, it is the actual shape of the problem: a decode step is
* memory-bound and tiny, and the fixed cost of dispatch is the thing you are
* fighting. It is the same argument as S13's launch overhead, arriving in a
* browser. `benchmark()` measures the crossover instead of asserting it.
*
* Everything is feature-detected. `navigator.gpu` is absent during SSR, absent
* in several shipping browsers, and absent behind flags. The engine runs
* identically without it.
*/
import type { QuantWeight } from "../model/llama.ts";
import { dequantizeRow } from "../format/quant.ts";
import { matmul } from "./matmul.ts";
/* ---------------------------------------------------------------- types */
/* WebGPU is not in every TS DOM lib yet, so declare the narrow surface we use
* rather than taking a dependency on @webgpu/types. */
type GPUBufferT = {
destroy(): void;
mapAsync(mode: number): Promise<void>;
getMappedRange(): ArrayBuffer;
unmap(): void;
};
type GPUPassT = {
setPipeline(p: unknown): void;
setBindGroup(i: number, g: unknown): void;
dispatchWorkgroups(x: number): void;
end(): void;
};
type GPUEncoderT = {
beginComputePass(): GPUPassT;
copyBufferToBuffer(a: GPUBufferT, ao: number, b: GPUBufferT, bo: number, n: number): void;
finish(): unknown;
};
type GPUPipelineT = { getBindGroupLayout(i: number): unknown };
type GPUDeviceT = {
createBuffer(d: { size: number; usage: number }): GPUBufferT;
createShaderModule(d: { code: string }): unknown;
createComputePipeline(d: unknown): GPUPipelineT;
createBindGroup(d: unknown): unknown;
createCommandEncoder(): GPUEncoderT;
queue: {
writeBuffer(b: GPUBufferT, o: number, d: ArrayBufferView): void;
submit(c: unknown[]): void;
};
destroy(): void;
};
type GPUAdapterT = { requestDevice(): Promise<GPUDeviceT>; info?: { vendor?: string } };
type NavigatorGPU = { gpu?: { requestAdapter(): Promise<GPUAdapterT | null> } };
/** Spec-defined GPUBufferUsage bits — hardcoded so no global type is needed. */
const USAGE = {
MAP_READ: 0x0001,
COPY_SRC: 0x0004,
COPY_DST: 0x0008,
UNIFORM: 0x0040,
STORAGE: 0x0080,
} as const;
const MAP_READ_MODE = 0x0001;
/* --------------------------------------------------------------- shader */
export const MATMUL_WGSL = /* wgsl */ `
struct Dims { n : u32, d : u32 };
@group(0) @binding(0) var<storage, read> W : array<f32>;
@group(0) @binding(1) var<storage, read> X : array<f32>;
@group(0) @binding(2) var<storage, read_write> Y : array<f32>;
@group(0) @binding(3) var<uniform> dims : Dims;
var<workgroup> partial : array<f32, 64>;
// One workgroup per output row; 64 lanes cooperate on one dot product.
@compute @workgroup_size(64)
fn main(@builtin(workgroup_id) wg : vec3<u32>,
@builtin(local_invocation_id) lid : vec3<u32>) {
let row = wg.x;
var acc = 0.0;
if (row < dims.d) {
var i = lid.x;
loop {
if (i >= dims.n) { break; }
acc = acc + W[row * dims.n + i] * X[i];
i = i + 64u;
}
}
partial[lid.x] = acc;
workgroupBarrier();
var stride = 32u;
loop {
if (stride == 0u) { break; }
if (lid.x < stride) { partial[lid.x] = partial[lid.x] + partial[lid.x + stride]; }
workgroupBarrier();
stride = stride / 2u;
}
if (lid.x == 0u && row < dims.d) { Y[row] = partial[0]; }
}
`;
/* ------------------------------------------------------------ detection */
export type GpuSupport =
| { supported: true; adapter: string }
| { supported: false; reason: string };
/** Cheap synchronous check — safe during render. */
export function hasWebGPU(): boolean {
return typeof navigator !== "undefined" && "gpu" in navigator;
}
export async function probeWebGPU(): Promise<GpuSupport> {
if (typeof navigator === "undefined") {
return { supported: false, reason: "no navigator (server-side render)" };
}
const gpu = (navigator as unknown as NavigatorGPU).gpu;
if (!gpu) {
return { supported: false, reason: "navigator.gpu is undefined — no WebGPU here" };
}
try {
const adapter = await gpu.requestAdapter();
if (!adapter) return { supported: false, reason: "requestAdapter() returned null" };
return { supported: true, adapter: adapter.info?.vendor || "unknown adapter" };
} catch (err) {
return { supported: false, reason: String(err) };
}
}
/**
* Dequantise a weight to f32 — what the GPU path uploads.
*
* Deliberately the *unfused* path from S08. Uploading once and reusing across
* tokens makes it the right trade for a small model, and the wrong one for a
* large model where 4-bit existed precisely so the weights would fit. A
* production WebGPU engine keeps them quantised on device and decodes in the
* shader.
*/
export function weightToFloat(w: QuantWeight, n: number, d: number): Float32Array {
return w.kind === "f32" ? w.data : dequantizeRow(w.raw, n * d, w.qtype);
}
/* -------------------------------------------------------------- backend */
/** A single uploaded weight matrix plus the buffers its dispatch needs. */
export class WebGPUMatmul {
private readonly device: GPUDeviceT;
private readonly pipeline: GPUPipelineT;
private readonly wBuf: GPUBufferT;
private readonly xBuf: GPUBufferT;
private readonly yBuf: GPUBufferT;
private readonly dimsBuf: GPUBufferT;
private readonly readBuf: GPUBufferT;
private readonly bindGroup: unknown;
readonly n: number;
readonly d: number;
readonly adapter: string;
private constructor(init: {
device: GPUDeviceT;
pipeline: GPUPipelineT;
wBuf: GPUBufferT;
xBuf: GPUBufferT;
yBuf: GPUBufferT;
dimsBuf: GPUBufferT;
readBuf: GPUBufferT;
bindGroup: unknown;
n: number;
d: number;
adapter: string;
}) {
this.device = init.device;
this.pipeline = init.pipeline;
this.wBuf = init.wBuf;
this.xBuf = init.xBuf;
this.yBuf = init.yBuf;
this.dimsBuf = init.dimsBuf;
this.readBuf = init.readBuf;
this.bindGroup = init.bindGroup;
this.n = init.n;
this.d = init.d;
this.adapter = init.adapter;
}
/** Upload `w` ([d, n], row-major) and prepare a reusable dispatch. */
static async create(w: Float32Array, n: number, d: number): Promise<WebGPUMatmul | null> {
const support = await probeWebGPU();
if (!support.supported) return null;
const gpu = (navigator as unknown as NavigatorGPU).gpu!;
const adapter = await gpu.requestAdapter();
if (!adapter) return null;
const device = await adapter.requestDevice();
if (w.length !== n * d) {
throw new RangeError(`weight has ${w.length} elements, expected ${n * d}`);
}
const shader = device.createShaderModule({ code: MATMUL_WGSL });
const pipeline = device.createComputePipeline({
layout: "auto",
compute: { module: shader, entryPoint: "main" },
});
const wBuf = device.createBuffer({
size: w.byteLength,
usage: USAGE.STORAGE | USAGE.COPY_DST,
});
device.queue.writeBuffer(wBuf, 0, w);
const xBuf = device.createBuffer({ size: n * 4, usage: USAGE.STORAGE | USAGE.COPY_DST });
const yBuf = device.createBuffer({ size: d * 4, usage: USAGE.STORAGE | USAGE.COPY_SRC });
const dimsBuf = device.createBuffer({ size: 8, usage: USAGE.UNIFORM | USAGE.COPY_DST });
device.queue.writeBuffer(dimsBuf, 0, new Uint32Array([n, d]));
const readBuf = device.createBuffer({ size: d * 4, usage: USAGE.COPY_DST | USAGE.MAP_READ });
const bindGroup = device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [
{ binding: 0, resource: { buffer: wBuf } },
{ binding: 1, resource: { buffer: xBuf } },
{ binding: 2, resource: { buffer: yBuf } },
{ binding: 3, resource: { buffer: dimsBuf } },
],
});
return new WebGPUMatmul({
device, pipeline, wBuf, xBuf, yBuf, dimsBuf, readBuf, bindGroup, n, d,
adapter: support.adapter,
});
}
/** out = W @ x. Async, because reading back from the GPU is. */
async run(x: Float32Array): Promise<Float32Array> {
const { device } = this;
device.queue.writeBuffer(this.xBuf, 0, x);
const enc = device.createCommandEncoder();
const pass = enc.beginComputePass();
pass.setPipeline(this.pipeline);
pass.setBindGroup(0, this.bindGroup);
pass.dispatchWorkgroups(this.d);
pass.end();
enc.copyBufferToBuffer(this.yBuf, 0, this.readBuf, 0, this.d * 4);
device.queue.submit([enc.finish()]);
await this.readBuf.mapAsync(MAP_READ_MODE);
const out = new Float32Array(this.readBuf.getMappedRange().slice(0));
this.readBuf.unmap();
return out;
}
destroy(): void {
for (const b of [this.wBuf, this.xBuf, this.yBuf, this.dimsBuf, this.readBuf]) b.destroy();
this.device.destroy();
}
}
export type BenchmarkResult = {
available: boolean;
reason?: string;
adapter?: string;
n: number;
d: number;
cpuMs: number;
gpuMs?: number;
maxAbsError?: number;
speedup?: number;
};
/**
* Run the same matmul on both paths and report the difference.
*
* This is what the chapter's simulator calls. It answers the only question that
* matters — is the GPU actually faster here, and does it agree with the CPU —
* rather than assuming either.
*/
export async function benchmark(n = 2048, d = 2048, iters = 8): Promise<BenchmarkResult> {
const w = new Float32Array(n * d);
const x = new Float32Array(n);
let seed = 12345;
const rnd = () => {
seed = (seed * 1103515245 + 12345) & 0x7fffffff;
return seed / 0x7fffffff - 0.5;
};
for (let i = 0; i < w.length; i++) w[i] = rnd();
for (let i = 0; i < x.length; i++) x[i] = rnd();
const cpuOut = new Float32Array(d);
const t0 = performance.now();
for (let i = 0; i < iters; i++) matmul(cpuOut, x, w, n, d);
const cpuMs = (performance.now() - t0) / iters;
const gpu = await WebGPUMatmul.create(w, n, d);
if (!gpu) {
const probe = await probeWebGPU();
return {
available: false,
reason: probe.supported ? "device creation failed" : probe.reason,
n, d, cpuMs,
};
}
try {
await gpu.run(x); // warm up: first dispatch includes pipeline compilation
const t1 = performance.now();
let gpuOut: Float32Array = new Float32Array(d);
for (let i = 0; i < iters; i++) gpuOut = await gpu.run(x);
const gpuMs = (performance.now() - t1) / iters;
let maxAbsError = 0;
for (let i = 0; i < d; i++) {
const e = Math.abs(gpuOut[i] - cpuOut[i]);
if (e > maxAbsError) maxAbsError = e;
}
return {
available: true,
adapter: gpu.adapter,
n, d, cpuMs, gpuMs, maxAbsError,
speedup: cpuMs / gpuMs,
};
} finally {
gpu.destroy();
}
}gpu probes for an adapter and, if it finds one, benchmarks the shader against the CPU path on your machine. If it does not, it says so and the engine carries on.
The GPU is a separate path, not a drop-in replacement
Workspace.matmul is synchronous and called seven times per layer. Routing decode through the GPU would mean awaiting dozens of round trips per token, and a round trip costs more than these matmuls save. So the GPU is an explicit async call for prefill-shaped work and large matrices, the CPU path stays the reference, and benchmark() measures the crossover on your machine instead of asserting it. This is S13's launch-overhead argument arriving in a browser.Step 10
Proving it works
Two files, and neither imports Node or the DOM. That is what lets the same assertions run headless and inside this page.
selftest.ts
90 assertions across every module: arena aliasing and exhaustion, matmul at awkward lengths against a naive loop, the streaming softmax against a materialised one, RoPE's relative-position property and its incompatibility with NeoX, every quantised format against hand-written bytes, the fused path against the unfused one, a GGUF round trip, cache immutability, tokenizer round trips, the top-p +1, nine grammar cases against JSON.parse, and greedy reproducibility.
/**
* mini-picoLM — the self-test
* ===========================
*
* Every claim the S22 chapter makes, checked. Same discipline as the Python
* chapters: a file that asserts its own claims rather than printing numbers and
* hoping.
*
* Runs in both places, unchanged:
*
* node --experimental-strip-types src/lib/minipicolm/selftest.ts
*
* and inside the browser, where the chapter's simulator renders the results.
* That is the point of keeping it free of DOM and Node APIs.
*/
import { Arena } from "./memory/arena.ts";
import { Tensor } from "./memory/tensor.ts";
import { alignUp } from "./memory/buffer.ts";
import { matmul, matmulQuant } from "./tensor/matmul.ts";
import { rmsnorm } from "./tensor/rmsnorm.ts";
import { softmax, OnlineSoftmax } from "./tensor/softmax.ts";
import { siluMulInto, siluScalar } from "./tensor/activation.ts";
import { RopeTables, ropeInPlace } from "./model/rope.ts";
import {
bitsPerWeight,
blockSize,
dequantizeRow,
fp16ToFp32,
fp32ToFp16,
quantSize,
type QuantType,
} from "./format/quant.ts";
import { GGUFFile } from "./format/gguf.ts";
import { JsonGrammar } from "./runtime/grammar.ts";
import { Sampler } from "./runtime/sampler.ts";
import { createToyModel } from "./main.ts";
import { generateText } from "./runtime/generation.ts";
export type Check = {
group: string;
name: string;
pass: boolean;
detail: string;
};
class Suite {
readonly checks: Check[] = [];
private group = "";
section(name: string): void {
this.group = name;
}
ok(name: string, pass: boolean, detail = ""): void {
this.checks.push({ group: this.group, name, pass, detail });
}
close(name: string, a: number, b: number, tol: number, detail?: string): void {
const d = Math.abs(a - b);
this.ok(name, d <= tol, detail ?? `|${a.toPrecision(6)} − ${b.toPrecision(6)}| = ${d.toExponential(2)}`);
}
get passed(): number {
return this.checks.filter((c) => c.pass).length;
}
get failed(): Check[] {
return this.checks.filter((c) => !c.pass);
}
}
/* ------------------------------------------------ hand-built quant blocks */
/** One Q4_K super-block with known fields, so the expected output is derivable. */
function handmadeQ4K(): { raw: Uint8Array; want: Float32Array } {
const sc = 3;
const mn = 1;
const raw = new Uint8Array(144);
const dv = new DataView(raw.buffer);
dv.setUint16(0, fp32ToFp16(0.5), true); // d
dv.setUint16(2, fp32ToFp16(0.25), true); // dmin
for (let j = 0; j < 4; j++) {
raw[4 + j] = sc;
raw[4 + j + 4] = mn;
}
for (let j = 4; j < 8; j++) {
raw[4 + j + 4] = (sc & 0xf) | ((mn & 0xf) << 4);
raw[4 + j - 4] |= ((sc >> 4) & 3) << 6;
raw[4 + j] |= ((mn >> 4) & 3) << 6;
}
raw.fill(0x21, 16, 144); // low nibble 1, high nibble 2
const lo = 0.5 * sc * 1 - 0.25 * mn;
const hi = 0.5 * sc * 2 - 0.25 * mn;
const want = new Float32Array(256);
for (let j = 0; j < 4; j++) {
want.fill(lo, j * 64, j * 64 + 32);
want.fill(hi, j * 64 + 32, j * 64 + 64);
}
return { raw, want };
}
function handmadeQ8_0(): { raw: Uint8Array; want: Float32Array } {
const raw = new Uint8Array(34);
const dv = new DataView(raw.buffer);
dv.setUint16(0, fp32ToFp16(0.125), true);
const want = new Float32Array(32);
for (let i = 0; i < 32; i++) {
const q = i - 16;
dv.setInt8(2 + i, q);
want[i] = 0.125 * q;
}
return { raw, want };
}
function handmadeQ4_0(): { raw: Uint8Array; want: Float32Array } {
const raw = new Uint8Array(18);
new DataView(raw.buffer).setUint16(0, fp32ToFp16(0.25), true);
raw.fill(0x30, 2, 18); // low nibble 0, high nibble 3
const want = new Float32Array(32);
want.fill(0.25 * (0 - 8), 0, 16);
want.fill(0.25 * (3 - 8), 16, 32);
return { raw, want };
}
/* ----------------------------------------------------------- the suite */
export async function runSelfTest(): Promise<Suite> {
const s = new Suite();
/* ---- steps 1-3: memory ---- */
s.section("memory");
{
const arena = new Arena(64 * 1024);
const a = arena.f32(100, "a");
const b = arena.f32(50, "b");
s.ok("one buffer backs every view", a.buffer === b.buffer && a.buffer === arena.buffer);
s.ok(
"allocations are aligned",
b.byteOffset % 256 === 0,
`b at byte ${b.byteOffset}`,
);
s.ok("views do not overlap", a.byteOffset + a.byteLength <= b.byteOffset);
a.fill(7);
b.fill(9);
s.ok("writing one view leaves the other alone", a[0] === 7 && a[99] === 7 && b[0] === 9);
const mark = arena.mark();
arena.f32(10, "scratch");
arena.release(mark);
s.ok("release rewinds the watermark", arena.used === mark);
let threw = false;
try {
arena.f32(1_000_000, "too big");
} catch {
threw = true;
}
s.ok("exhaustion throws rather than corrupting", threw);
s.ok("alignUp rounds up", alignUp(1, 256) === 256 && alignUp(256, 256) === 256);
const t = Tensor.zeros(arena, [4, 8], "t");
t.row(2).fill(3);
s.ok("tensor rows are views", t.data[16] === 3 && t.data[15] === 0, "row(2) starts at 16");
let shapeThrew = false;
try {
Tensor.of(new Float32Array(7), [4, 2]);
} catch {
shapeThrew = true;
}
s.ok("shape mismatch is caught", shapeThrew);
}
/* ---- steps 4-5, 9: kernels ---- */
s.section("kernels");
{
// matmul against a hand-computed 2x3 case
const w = new Float32Array([1, 2, 3, 4, 5, 6]); // [2,3]
const x = new Float32Array([1, 0, -1]);
const out = new Float32Array(2);
matmul(out, x, w, 3, 2);
s.ok("matmul is correct", out[0] === -2 && out[1] === -2, `got [${out[0]}, ${out[1]}]`);
// unrolling must not change the answer at awkward lengths
for (const n of [1, 3, 5, 7, 16, 17, 33]) {
const ww = new Float32Array(n * 2);
const xx = new Float32Array(n);
for (let i = 0; i < n; i++) {
xx[i] = Math.sin(i);
ww[i] = Math.cos(i);
ww[n + i] = Math.cos(i * 2);
}
const o = new Float32Array(2);
matmul(o, xx, ww, n, 2);
let ref0 = 0;
let ref1 = 0;
for (let i = 0; i < n; i++) {
ref0 += ww[i] * xx[i];
ref1 += ww[n + i] * xx[i];
}
s.close(`matmul n=${n} matches a naive loop`, o[0], ref0, 1e-5);
s.close(`matmul n=${n} row 1`, o[1], ref1, 1e-5);
}
// rmsnorm: an all-ones vector normalises to ~1
const v = new Float32Array(64).fill(1);
const g = new Float32Array(64).fill(1);
const no = new Float32Array(64);
rmsnorm(no, v, g);
s.close("rmsnorm of a constant vector is ~1", no[0], 1, 1e-4);
const sm = new Float32Array([1, 2, 3]);
softmax(sm);
let sum = 0;
for (const p of sm) sum += p;
s.close("softmax sums to 1", sum, 1, 1e-6);
s.ok("softmax is monotone", sm[0] < sm[1] && sm[1] < sm[2]);
const big = new Float32Array([1000, 1001, 1002]);
softmax(big);
s.ok("softmax does not overflow at large logits", Number.isFinite(big[0]) && big[2] > big[0]);
// SwiGLU
const hb = new Float32Array([1, -1]);
const hb2 = new Float32Array([2, 2]);
siluMulInto(hb, hb2);
s.close("SwiGLU gate", hb[0], siluScalar(1) * 2, 1e-6);
s.close("SwiGLU gate (negative)", hb[1], siluScalar(-1) * 2, 1e-6);
}
/* ---- online softmax: the S12 claim ---- */
s.section("online softmax");
{
const n = 40;
const dim = 8;
const scores = new Float32Array(n);
const values = new Float32Array(n * dim);
for (let i = 0; i < n; i++) {
scores[i] = Math.sin(i) * 5;
for (let j = 0; j < dim; j++) values[i * dim + j] = Math.cos(i * 0.3 + j);
}
// reference: materialise the whole score row, then weight
const probs = scores.slice();
softmax(probs);
const ref = new Float32Array(dim);
for (let i = 0; i < n; i++) {
for (let j = 0; j < dim; j++) ref[j] += probs[i] * values[i * dim + j];
}
const online = new OnlineSoftmax(dim, new Float32Array(dim));
online.reset();
for (let i = 0; i < n; i++) online.push(scores[i], values, i * dim);
const got = new Float32Array(dim);
online.finish(got, 0);
let worst = 0;
for (let j = 0; j < dim; j++) worst = Math.max(worst, Math.abs(got[j] - ref[j]));
s.ok(
"streaming softmax equals the materialised one",
worst < 1e-6,
`max abs error ${worst.toExponential(2)} over ${n} positions`,
);
// ascending scores are the worst case for the rescaling path
const asc = new OnlineSoftmax(dim, new Float32Array(dim));
asc.reset();
for (let i = 0; i < n; i++) asc.push(i, values, (i % n) * dim);
const ascOut = new Float32Array(dim);
asc.finish(ascOut, 0);
s.ok("monotonically rising scores stay finite", ascOut.every((v) => Number.isFinite(v)));
// a fully-masked stream — every score -Infinity — must come out zero, not NaN
const masked = new OnlineSoftmax(dim, new Float32Array(dim));
masked.reset();
for (let i = 0; i < 4; i++) masked.push(-Infinity, values, 0);
const maskedOut = new Float32Array(dim);
masked.finish(maskedOut, 0);
s.ok("a fully-masked stream yields zeros, not NaN", maskedOut.every((v) => v === 0));
}
/* ---- step 6: RoPE ---- */
s.section("rope");
{
const headDim = 8;
const tables = new RopeTables(64, headDim, 10000);
const a = new Float32Array(headDim);
const b = new Float32Array(headDim);
for (let i = 0; i < headDim; i++) {
a[i] = Math.sin(i + 1);
b[i] = Math.cos(i * 2 + 1);
}
// the property that makes a KV cache sound
const scores: number[] = [];
for (const base of [0, 5, 20, 50]) {
const q = a.slice();
const k = b.slice();
ropeInPlace(q, 1, headDim, tables, base);
ropeInPlace(k, 1, headDim, tables, base + 3);
let dot = 0;
for (let i = 0; i < headDim; i++) dot += q[i] * k[i];
scores.push(dot);
}
const spread = Math.max(...scores) - Math.min(...scores);
s.ok(
"the score depends on the gap, not the offset",
spread < 1e-5,
`spread ${spread.toExponential(2)} across offsets 0..50 at gap 3`,
);
// interleaved is NOT the NeoX layout S03 teaches
const inter = a.slice();
ropeInPlace(inter, 1, headDim, tables, 3);
const half = headDim >> 1;
const neox = new Float32Array(headDim);
for (let i = 0; i < half; i++) {
const c = tables.cos[3 * half + i];
const sn = tables.sin[3 * half + i];
neox[i] = a[i] * c - a[i + half] * sn;
neox[i + half] = a[i + half] * c + a[i] * sn;
}
let diff = 0;
for (let i = 0; i < headDim; i++) diff = Math.max(diff, Math.abs(inter[i] - neox[i]));
s.ok(
"interleaved differs from NeoX",
diff > 0.1,
`max abs difference ${diff.toFixed(4)} — pick the wrong one and the model is subtly broken`,
);
// position 0 is the identity rotation
const p0 = a.slice();
ropeInPlace(p0, 1, headDim, tables, 0);
let same = 0;
for (let i = 0; i < headDim; i++) same = Math.max(same, Math.abs(p0[i] - a[i]));
s.ok("position 0 rotates by nothing", same < 1e-6);
}
/* ---- steps 15-16: quantisation ---- */
s.section("quantisation");
{
const cases: [QuantType, { raw: Uint8Array; want: Float32Array }][] = [
["Q4_K", handmadeQ4K()],
["Q8_0", handmadeQ8_0()],
["Q4_0", handmadeQ4_0()],
];
for (const [t, { raw, want }] of cases) {
const got = dequantizeRow(raw, blockSize(t), t);
let worst = 0;
for (let i = 0; i < want.length; i++) worst = Math.max(worst, Math.abs(got[i] - want[i]));
s.ok(
`${t} matches bytes written by hand from the spec`,
worst < 1e-6,
`${quantSize(t)} B → ${blockSize(t)} weights, max err ${worst.toExponential(2)}`,
);
}
s.close("Q4_K is exactly 4.5 bits per weight", bitsPerWeight("Q4_K"), 4.5, 1e-9);
s.close("Q6_K is 6.5625 bits per weight", bitsPerWeight("Q6_K"), 6.5625, 1e-9);
// fp16 round trip
for (const v of [0, 1, -2.5, 65504, 6.1e-5]) {
const back = fp16ToFp32(fp32ToFp16(v));
s.close(`fp16 round-trips ${v}`, back, v, Math.max(1e-7, Math.abs(v) * 1e-3));
}
s.ok("1e-8 flushes to zero (below the subnormal floor)", fp16ToFp32(fp32ToFp16(1e-8)) === 0);
s.ok("65504 is the largest finite fp16", fp16ToFp32(fp32ToFp16(65504)) === 65504);
// the fused path must equal the unfused one
const n = 256;
const d = 3;
const q4k = handmadeQ4K();
const raw = new Uint8Array(144 * d);
for (let i = 0; i < d; i++) raw.set(q4k.raw, i * 144);
const x = new Float32Array(n);
for (let i = 0; i < n; i++) x[i] = Math.sin(i);
const fused = new Float32Array(d);
matmulQuant(fused, x, raw, n, d, "Q4_K", new Float32Array(256));
const dequantised = dequantizeRow(raw, n * d, "Q4_K");
const unfused = new Float32Array(d);
matmul(unfused, x, dequantised, n, d);
let worst = 0;
for (let i = 0; i < d; i++) worst = Math.max(worst, Math.abs(fused[i] - unfused[i]));
s.ok(
"fused dequant+matmul equals dequantise-then-matmul",
worst < 1e-3,
`max abs difference ${worst.toExponential(2)} — same maths, different bytes moved`,
);
}
/* ---- step 14: GGUF ---- */
s.section("gguf");
{
const engine = createToyModel({ seed: 3 });
// Re-read the very bytes createToyModel wrote, through the real reader.
const f = new GGUFFile(engine.ggufBuffer);
s.ok("header parses", f.metadata.get("general.architecture") === "llama");
s.ok("alignment is honoured", f.dataStart % 32 === 0, `data starts at byte ${f.dataStart}`);
s.ok("tensor index is complete", f.tensors.size === 3 + 9 * 2, `${f.tensors.size} tensors`);
s.ok(
"every tensor lies inside the file",
[...f.tensors.values()].every(
(t) => f.dataStart + t.offset + t.nBytes <= engine.ggufBuffer.byteLength,
),
);
const embd = f.info("token_embd.weight");
s.ok(
"shapes survive the dims reversal",
embd.dims.length === 2 && embd.nElements === engine.config.vocabSize * engine.config.dim,
`dims [${embd.dims}]`,
);
s.ok(
"config was read back",
engine.config.nLayers === 2 && engine.config.dim === 64 && engine.config.nKvHeads === 2,
`${engine.config.nLayers}L d=${engine.config.dim}`,
);
s.ok(
"the vocabulary survived the metadata round trip",
(f.metadata.get("tokenizer.ggml.tokens") as string[]).length === engine.tokenizer.vocabSize,
);
let badMagic = false;
try {
new GGUFFile(new ArrayBuffer(64));
} catch {
badMagic = true;
}
s.ok("a non-GGUF buffer is rejected with a clear error", badMagic);
}
/* ---- steps 7, 8, 10: the forward pass ---- */
s.section("forward pass");
{
const { model, tokenizer } = createToyModel({ seed: 5 });
const ids = tokenizer.encode("the cat sat", true);
s.ok("BOS is prepended", ids[0] === tokenizer.bosId);
let logits = model.forward(ids[0], 0);
for (let i = 1; i < ids.length; i++) logits = model.forward(ids[i], i);
s.ok("logits have the vocab shape", logits.length === model.cfg.vocabSize);
s.ok("activations stay finite", logits.every((v) => Number.isFinite(v)));
const first = logits.slice();
// replaying the same prefix from a clean cache must be bit-identical
model.reset();
let again = model.forward(ids[0], 0);
for (let i = 1; i < ids.length; i++) again = model.forward(ids[i], i);
let worst = 0;
for (let i = 0; i < first.length; i++) worst = Math.max(worst, Math.abs(first[i] - again[i]));
s.ok(
"replaying a prefix is bit-identical",
worst === 0,
`max abs difference ${worst} — the cache is an optimisation, not an approximation`,
);
// history is immutable
const snapshot = model.cache.key.slice(0, model.cfg.kvDim * ids.length);
model.forward(5, ids.length);
let changed = false;
for (let i = 0; i < snapshot.length; i++) {
if (model.cache.key[i] !== snapshot[i]) changed = true;
}
s.ok("writing position N leaves earlier positions alone", !changed);
// GQA actually shares
s.ok(
"GQA divides the cache",
model.cfg.nHeads % model.cfg.nKvHeads === 0 && model.cfg.nKvHeads < model.cfg.nHeads,
`${model.cfg.nHeads} query heads over ${model.cfg.nKvHeads} kv heads = ${model.cfg.nHeads / model.cfg.nKvHeads}×`,
);
}
/* ---- steps 11-12: tokenizer and sampler ---- */
s.section("tokenizer & sampler");
{
const { tokenizer } = createToyModel({ seed: 7 });
const text = "the cat sat on the mat";
const ids = tokenizer.encode(text, true);
s.ok("round-trips", tokenizer.decode(ids.slice(1)) === text, JSON.stringify(tokenizer.decode(ids.slice(1))));
const weird = tokenizer.encode("q~z", false);
s.ok(
"byte fallback makes the vocabulary total",
tokenizer.decode(weird, false).endsWith("q~z"),
JSON.stringify(tokenizer.decode(weird, false)),
);
const sampler = new Sampler({ temperature: 0, seed: 1 });
const logits = new Float32Array(64);
logits[13] = 5;
const draws = new Set([0, 1, 2, 3, 4].map(() => sampler.sample(logits)));
s.ok("greedy ignores the seed", draws.size === 1 && draws.has(13));
// the +1 that keeps the crossing token
const peaked = new Float32Array(64).fill(-20);
peaked[7] = 20;
const nucleus = new Sampler({ temperature: 1, topP: 0.9, seed: 2 });
s.ok(
"top-p keeps the token that crosses the threshold",
nucleus.reachable(peaked) === 1 && nucleus.sample(peaked) === 7,
`reachable = ${nucleus.reachable(peaked)}`,
);
// seeded reproducibility
const a = new Sampler({ temperature: 1, topP: 1, seed: 42 });
const b = new Sampler({ temperature: 1, topP: 1, seed: 42 });
const flat = new Float32Array(32);
for (let i = 0; i < 32; i++) flat[i] = Math.sin(i);
const seqA = [0, 1, 2, 3, 4].map(() => a.sample(flat));
const seqB = [0, 1, 2, 3, 4].map(() => b.sample(flat));
s.ok("the same seed gives the same sequence", seqA.join() === seqB.join(), seqA.join(","));
}
/* ---- step 17: grammar ---- */
s.section("grammar");
{
const g = new JsonGrammar();
const cases: [string, boolean][] = [
['{"a":1}', true],
['{"a":1', false],
["[1,2]", true],
["{]", false],
["true", true],
['{"a":}', false],
['{"a":{"b":[1,true]}}', true],
["1.5e3", true],
["", false],
];
for (const [text, want] of cases) {
const end = g.advance(g.start(), text);
const got = g.accepting(end);
s.ok(`grammar ${want ? "accepts" : "rejects"} ${JSON.stringify(text)}`, got === want);
if (want) {
let parses = true;
try {
JSON.parse(text);
} catch {
parses = false;
}
s.ok(`JSON.parse agrees on ${JSON.stringify(text)}`, parses);
}
}
const pieces = ["{", "}", '"', ":", ",", "a", "1", "[", "]", "true"];
const mask = new Float32Array(pieces.length).fill(0);
const legal = g.maskLogits(mask, g.start(), pieces);
s.ok("'{' is legal at the start", mask[0] === 0);
s.ok("'}' is not", mask[1] === -Infinity);
s.ok("some tokens survive the mask", legal > 0 && legal < pieces.length, `${legal}/${pieces.length}`);
}
/* ---- step 13: generation ---- */
s.section("generation");
{
const { model, tokenizer, sampler } = createToyModel({ seed: 11 });
sampler.temperature = 0;
const r = await generateText(model, tokenizer, sampler, {
prompt: "the cat",
maxTokens: 12,
yieldEvery: 4,
});
s.ok("generation produces tokens", r.outputTokens > 0, `${r.outputTokens} tokens`);
s.ok("prompt was tokenized", r.promptTokens > 1, `${r.promptTokens} prompt tokens`);
s.ok(
"a stop reason is always reported",
["eos", "maxTokens", "contextFull", "aborted", "grammarComplete"].includes(r.stopReason),
r.stopReason,
);
// determinism at temperature 0
const second = createToyModel({ seed: 11 });
second.sampler.temperature = 0;
const r2 = await generateText(second.model, second.tokenizer, second.sampler, {
prompt: "the cat",
maxTokens: 12,
yieldEvery: 4,
});
s.ok("greedy generation is reproducible", r.text === r2.text, JSON.stringify(r.text.slice(0, 40)));
// constrained decoding yields parseable JSON or an explicit incompletion
const j = createToyModel({ seed: 13 });
j.sampler.temperature = 0;
const jr = await generateText(j.model, j.tokenizer, j.sampler, {
prompt: "json",
maxTokens: 40,
json: true,
});
if (jr.grammarValid) {
let parses = true;
try {
JSON.parse(jr.text);
} catch {
parses = false;
}
s.ok("completed constrained output parses", parses, JSON.stringify(jr.text));
} else {
s.ok(
"incomplete constrained output is reported, not returned as valid",
jr.grammarValid === false,
`stopped: ${jr.stopReason} — masking guarantees valid, not complete`,
);
}
}
return s;
}
/* --------------------------------------------------------------- runner */
export type SelfTestSummary = {
total: number;
passed: number;
failed: Check[];
checks: Check[];
};
export async function selfTest(): Promise<SelfTestSummary> {
const s = await runSelfTest();
return { total: s.checks.length, passed: s.passed, failed: s.failed, checks: s.checks };
}run-selftest.mts
The only file that knows about a terminal.
/**
* Headless runner for the mini-picoLM self-test.
*
* node --experimental-strip-types src/lib/minipicolm/run-selftest.ts
*
* The engine itself uses no Node and no DOM APIs, so the exact same checks run
* in the browser behind the S22 simulator. This file is the only part that
* knows about a terminal.
*/
import { selfTest } from "./selftest.ts";
const GREEN = "[32m";
const RED = "[31m";
const DIM = "[2m";
const BOLD = "[1m";
const RESET = "[0m";
const t0 = performance.now();
const { total, passed, failed, checks } = await selfTest();
const ms = performance.now() - t0;
let group = "";
for (const c of checks) {
if (c.group !== group) {
group = c.group;
process.stdout.write(`\n${BOLD}▸ ${group}${RESET}\n`);
}
const mark = c.pass ? `${GREEN}✓${RESET}` : `${RED}✗${RESET}`;
const detail = c.detail ? ` ${DIM}${c.detail}${RESET}` : "";
process.stdout.write(` ${mark} ${c.name}${detail}\n`);
}
process.stdout.write(
`\n ${BOLD}${passed}/${total} checks passed${RESET} in ${ms.toFixed(0)} ms\n`,
);
if (failed.length) {
process.stdout.write(`\n${RED}${BOLD}failures:${RESET}\n`);
for (const f of failed) {
process.stdout.write(` ${RED}✗${RESET} ${f.group} / ${f.name}${f.detail ? ` — ${f.detail}` : ""}\n`);
}
process.exit(1);
}# 90 assertions across every module
node --experimental-strip-types src/lib/minipicolm/run-selftest.mts
# every terminal command, dispatched and checked
node --experimental-strip-types src/lib/minipicolm/run-shell-test.mts$ node --experimental-strip-types src/lib/minipicolm/run-selftest.mtsOnly Node 22.6+ is required: no npm install, no build step.
Step 11
The terminal
Here is the finished engine, loaded in this page, behind a shell. Every command calls real code: selftest runs the same 90 assertions the headless runner does, quant decodes an actual block, gen streams from the real forward pass, src prints the module you just read.
mini-picoLM · the engine is loaded in this tab
type help · tab completes · ↑ recalls · every command runs real code
Use it to check your own copy. If you have been building along, run selftest here and then against your files; the groups and counts should match. bench will disagree, because that is your machine.
The shell is a module too
shell.ts holds the command table and knows nothing about the DOM, so run-shell-test.mts can drive all twelve commands headlessly. A verification surface that cannot itself be verified is decoration.shell.ts
The command table. Every entry calls into a module you have already read, and nothing here is a canned response, so gguf can fail if the writer breaks.
/**
* mini-picoLM — a shell over the engine
* =====================================
*
* The command layer behind the terminal at the end of S22. It is deliberately
* separate from the React component: this file knows about the engine and
* nothing about the DOM, so the same commands can be driven from a test.
*
* Every command exercises a real module. `selftest` runs the same 89 assertions
* the headless runner does; `quant` decodes an actual GGUF block; `gen` streams
* from the real forward pass. Nothing here is a canned response.
*/
import { createToyModel, type Engine } from "./main.ts";
import { generate } from "./runtime/generation.ts";
import { selfTest } from "./selftest.ts";
import { GGUFFile } from "./format/gguf.ts";
import {
bitsPerWeight,
blockSize,
dequantizeRow,
fp16ToFp32,
fp32ToFp16,
quantSize,
type QuantType,
} from "./format/quant.ts";
import { benchmark, probeWebGPU } from "./tensor/webgpu.ts";
export type Line = { text: string; tone?: "ok" | "warn" | "err" | "dim" | "accent" };
export type Emit = (line: Line) => void;
/** What `src` prints. Injected so the engine never imports chapter content. */
export type SourceIndex = Record<string, { path: string; lines: number; text: string }>;
export type ShellState = {
engine: Engine & { ggufBytes: number; ggufBuffer: ArrayBuffer };
source: SourceIndex;
sourceLines: number;
};
const bytes = (n: number) =>
n < 1024 ? `${n} B` : n < 1024 ** 2 ? `${(n / 1024).toFixed(1)} KB` : `${(n / 1024 ** 2).toFixed(2)} MB`;
const pad = (s: string, n: number) => s.padEnd(n, " ");
export type Command = {
name: string;
args?: string;
help: string;
run: (argv: string[], out: Emit, state: ShellState) => Promise<void> | void;
};
/* ------------------------------------------------------------- commands */
export const COMMANDS: Command[] = [
{
name: "help",
help: "list every command",
run: (_a, out) => {
out({ text: "commands — everything here calls the real engine", tone: "dim" });
for (const c of COMMANDS) {
out({ text: ` ${pad(c.name + (c.args ? " " + c.args : ""), 26)}${c.help}` });
}
},
},
{
name: "selftest",
args: "[group]",
help: "run the 89 assertions in this tab",
run: async (argv, out) => {
const filter = argv[0]?.toLowerCase();
out({ text: "running…", tone: "dim" });
const r = await selfTest();
let group = "";
let shown = 0;
for (const c of r.checks) {
if (filter && !c.group.toLowerCase().includes(filter)) continue;
if (c.group !== group) {
group = c.group;
out({ text: `▸ ${group}`, tone: "accent" });
}
shown++;
out({
text: ` ${c.pass ? "✓" : "✗"} ${c.name}${c.detail ? ` ${c.detail}` : ""}`,
tone: c.pass ? "ok" : "err",
});
}
if (filter && shown === 0) out({ text: `no group matching "${filter}"`, tone: "warn" });
out({
text: `${r.passed}/${r.total} checks passed`,
tone: r.passed === r.total ? "ok" : "err",
});
},
},
{
name: "gen",
args: '"prompt" [-n 24] [--json] [-t 0.9]',
help: "stream tokens from the forward pass",
run: async (argv, out, state) => {
const joined = argv.join(" ");
const quoted = /"([^"]*)"/.exec(joined);
const prompt = quoted ? quoted[1] : "the cat sat";
const n = Number(/-n\s+(\d+)/.exec(joined)?.[1] ?? 24);
const temp = Number(/-t\s+([\d.]+)/.exec(joined)?.[1] ?? 0.9);
const json = joined.includes("--json");
const { model, tokenizer, sampler } = state.engine;
model.reset();
sampler.temperature = temp;
sampler.reseed(7);
out({ text: `prompt ${JSON.stringify(prompt)} · n=${n} · t=${temp}${json ? " · json" : ""}`, tone: "dim" });
const it = generate(model, tokenizer, sampler, { prompt, maxTokens: n, json, yieldEvery: 4 });
let acc = "";
for (;;) {
const step = await it.next();
if (step.done) {
out({ text: JSON.stringify(acc), tone: "accent" });
out({
text:
`${step.value.outputTokens} tokens · ${step.value.tokensPerSecond.toFixed(1)} tok/s · ` +
`prefill ${step.value.prefillMs.toFixed(1)} ms · stopped: ${step.value.stopReason}` +
(step.value.grammarValid !== undefined
? ` · grammar ${step.value.grammarValid ? "complete" : "incomplete"}`
: ""),
tone: "dim",
});
if (json && step.value.grammarValid) {
try {
JSON.parse(acc);
out({ text: "JSON.parse accepted it", tone: "ok" });
} catch {
out({ text: "JSON.parse REJECTED it — that would be a grammar bug", tone: "err" });
}
}
break;
}
acc += step.value.text;
}
},
},
{
name: "tokenize",
args: "<text>",
help: "SentencePiece BPE, with a round-trip check",
run: (argv, out, state) => {
const text = argv.join(" ") || "the cat sat on the mat";
const tok = state.engine.tokenizer;
const ids = tok.encode(text, true);
out({ text: `ids ${JSON.stringify(ids)}` });
out({ text: `pieces ${ids.map((i) => JSON.stringify(tok.piece(i))).join(" ")}` });
const back = tok.decode(ids.slice(1));
const same = back === text;
out({ text: `decode ${JSON.stringify(back)}`, tone: same ? "ok" : "err" });
out({
text: same ? "round-trip exact" : "round-trip DIFFERS — that is a bug",
tone: same ? "ok" : "err",
});
},
},
{
name: "quant",
args: "[Q4_0|Q8_0|Q4_K|Q6_K]",
help: "block layout, bits/weight, and a decoded block",
run: (argv, out) => {
const types: QuantType[] = ["Q4_0", "Q8_0", "Q4_K", "Q6_K"];
const want = argv[0]?.toUpperCase() as QuantType | undefined;
if (!want) {
out({ text: `${pad("type", 8)}${pad("bytes", 8)}${pad("weights", 9)}${pad("bpw", 8)}vs fp16` });
for (const t of types) {
out({
text:
pad(t, 8) +
pad(String(quantSize(t)), 8) +
pad(String(blockSize(t)), 9) +
pad(bitsPerWeight(t).toFixed(2), 8) +
`${(16 / bitsPerWeight(t)).toFixed(2)}×`,
});
}
out({ text: 'pass a type to decode one block, e.g. "quant Q4_K"', tone: "dim" });
return;
}
if (!types.includes(want)) {
out({ text: `unknown type ${want}`, tone: "err" });
return;
}
// A block whose fields are set so the expected values are derivable.
const size = quantSize(want);
const raw = new Uint8Array(size);
const dv = new DataView(raw.buffer);
dv.setUint16(0, fp32ToFp16(0.5), true);
if (want === "Q4_K") dv.setUint16(2, fp32ToFp16(0.25), true);
if (want === "Q6_K") dv.setUint16(208, fp32ToFp16(0.5), true);
for (let i = want === "Q6_K" ? 0 : 2; i < size; i++) raw[i] = 0x21;
const decoded = dequantizeRow(raw, blockSize(want), want);
out({ text: `${want}: ${size} bytes → ${blockSize(want)} weights (${bitsPerWeight(want).toFixed(2)} bpw)` });
out({
text: `first 12: ${Array.from(decoded.slice(0, 12)).map((v) => v.toFixed(3)).join(" ")}`,
tone: "accent",
});
},
},
{
name: "fp16",
args: "<number>",
help: "round-trip a value through IEEE-754 binary16",
run: (argv, out) => {
const v = Number(argv[0] ?? 1);
if (!Number.isFinite(v)) {
out({ text: "give me a finite number", tone: "err" });
return;
}
const bits = fp32ToFp16(v);
const back = fp16ToFp32(bits);
out({ text: `in ${v}` });
out({ text: `bits 0x${bits.toString(16).padStart(4, "0")}` });
out({ text: `out ${back}`, tone: back === v ? "ok" : "warn" });
if (back !== v) {
out({
text:
Math.abs(v) > 65504
? "outside fp16's range — saturated"
: Math.abs(v) < 6.1e-5
? "below the subnormal floor — flushed toward zero"
: `relative error ${(Math.abs(back - v) / Math.abs(v)).toExponential(2)}`,
tone: "dim",
});
}
},
},
{
name: "mem",
help: "the arena: every allocation, in order",
run: (_a, out, state) => {
const { model, ggufBytes } = state.engine;
const arena = model.arena;
out({ text: `arena ${bytes(arena.used)} used of ${bytes(arena.capacity)}` });
out({ text: `kv cache ${bytes(model.kvBytes)}` });
out({ text: `gguf ${bytes(ggufBytes)} (written in memory, then re-read)` });
out({ text: `allocations (${arena.allocations.length}):`, tone: "dim" });
for (const a of arena.allocations) {
out({ text: ` ${pad(a.name, 16)}@ ${pad(String(a.byteOffset), 10)}${bytes(a.byteLength)}` });
}
},
},
{
name: "model",
help: "the config read back out of the GGUF header",
run: (_a, out, state) => {
const c = state.engine.config;
out({ text: `layers ${c.nLayers}` });
out({ text: `d_model ${c.dim}` });
out({ text: `heads ${c.nHeads} query / ${c.nKvHeads} kv (GQA ${c.nHeads / c.nKvHeads}×)` });
out({ text: `head_dim ${c.headDim}` });
out({ text: `ffn hidden ${c.hidden}` });
out({ text: `vocab ${c.vocabSize}` });
out({ text: `context ${c.maxSeq}` });
out({ text: `rope base ${c.ropeFreqBase}` });
out({ text: state.engine.source, tone: "dim" });
},
},
{
name: "gguf",
help: "parse the in-memory file and list its tensors",
run: (_a, out, state) => {
const f = new GGUFFile(state.engine.ggufBuffer);
out({ text: `magic ok · ${f.tensors.size} tensors · ${f.metadata.size} metadata keys` });
out({ text: `data starts at byte ${f.dataStart} (aligned)`, tone: "dim" });
let n = 0;
for (const t of f.tensors.values()) {
if (n++ >= 8) {
out({ text: ` … ${f.tensors.size - 8} more`, tone: "dim" });
break;
}
out({ text: ` ${pad(t.name, 26)}${pad(`[${t.dims}]`, 14)}${pad(t.qtype, 6)}${bytes(t.nBytes)}` });
}
},
},
{
name: "gpu",
help: "WebGPU adapter, and a CPU-vs-GPU matmul benchmark",
run: async (_a, out) => {
const probe = await probeWebGPU();
if (!probe.supported) {
out({ text: `WebGPU unavailable — ${probe.reason}`, tone: "warn" });
out({ text: "the engine runs identically without it; that is the point", tone: "dim" });
return;
}
out({ text: `adapter ${probe.adapter}`, tone: "ok" });
out({ text: "benchmarking 1024×1024…", tone: "dim" });
const r = await benchmark(1024, 1024, 4);
if (!r.available) {
out({ text: `benchmark unavailable — ${r.reason}`, tone: "warn" });
return;
}
out({ text: `cpu ${r.cpuMs?.toFixed(2)} ms` });
out({ text: `gpu ${r.gpuMs?.toFixed(2)} ms` });
out({
text: `speedup ${(r.speedup ?? 0).toFixed(2)}× · max |gpu − cpu| ${(r.maxAbsError ?? 0).toExponential(1)}`,
tone: (r.speedup ?? 0) > 1 ? "ok" : "warn",
});
if ((r.speedup ?? 0) <= 1) {
out({ text: "the CPU won — upload and readback cost more than the matmul saves", tone: "dim" });
}
},
},
{
name: "src",
args: "[module]",
help: "the source of any module, as shipped",
run: (argv, out, state) => {
const key = argv[0];
const keys = Object.keys(state.source);
if (!key) {
out({ text: `${state.sourceLines} lines across ${keys.length} files`, tone: "dim" });
for (const k of keys) {
const m = state.source[k];
out({ text: ` ${pad(k, 14)}${pad(m.path, 30)}${String(m.lines).padStart(4)} lines` });
}
out({ text: 'e.g. "src arena"', tone: "dim" });
return;
}
const mod = state.source[key];
if (!mod) {
out({ text: `no module "${key}" — try "src" for the list`, tone: "err" });
return;
}
out({ text: `${mod.path} · ${mod.lines} lines`, tone: "accent" });
for (const line of mod.text.split("\n")) out({ text: line });
},
},
{
name: "bench",
args: "[tokens]",
help: "time the forward pass",
run: async (argv, out, state) => {
const n = Math.min(Number(argv[0] ?? 16), 64);
const { model } = state.engine;
model.reset();
model.ws.resetCounters();
const t0 = performance.now();
for (let i = 0; i < n; i++) model.forward(7, i);
const ms = performance.now() - t0;
out({ text: `${n} forward passes in ${ms.toFixed(1)} ms` });
out({ text: `${(ms / n).toFixed(2)} ms/token · ${((n / ms) * 1000).toFixed(1)} tok/s`, tone: "accent" });
out({
text: `${model.ws.matmulCalls} matmuls · ${(model.ws.matmulMacs / 1e6).toFixed(2)}M MACs`,
tone: "dim",
});
},
},
];
export const COMMAND_NAMES = COMMANDS.map((c) => c.name);
export function createShellState(source: SourceIndex = {}, sourceLines = 0): ShellState {
return {
engine: createToyModel({ nLayers: 2, dim: 64, hidden: 128, maxSeq: 128, seed: 1 }),
source,
sourceLines,
};
}
/** Dispatch one line. Unknown commands suggest the closest match. */
export async function runCommand(line: string, out: Emit, state: ShellState): Promise<void> {
const parts = line.trim().split(/\s+/).filter(Boolean);
if (!parts.length) return;
const [name, ...argv] = parts;
const cmd = COMMANDS.find((c) => c.name === name);
if (!cmd) {
const near = COMMAND_NAMES.filter((n) => n.startsWith(name[0]));
out({
text: `unknown command: ${name}${near.length ? ` — did you mean ${near.join(", ")}?` : ""}`,
tone: "err",
});
out({ text: 'type "help" for the list', tone: "dim" });
return;
}
try {
await cmd.run(argv, out, state);
} catch (err) {
out({ text: `error: ${err instanceof Error ? err.message : String(err)}`, tone: "err" });
}
}run-shell-test.mts
Each command dispatched, its output captured, and asserted against what it is supposed to prove.
/**
* Headless exercise of every browser-shell command.
*
* node --experimental-strip-types src/lib/minipicolm/run-shell-test.mts
*
* The terminal at the end of S22 is a real verification surface, so it gets a
* real gate: each command is dispatched, its output captured, and asserted
* against what it is supposed to prove. A command that silently stops working
* fails here as well as in the page.
*/
import { COMMANDS, createShellState, runCommand, type Line } from "./shell.ts";
const GREEN = "[32m";
const RED = "[31m";
const DIM = "[2m";
const BOLD = "[1m";
const RESET = "[0m";
const state = createShellState(
{ arena: { path: "memory/arena.ts", lines: 116, text: "// stub for the headless run" } },
116,
);
let passed = 0;
const failures: string[] = [];
async function expect(cmd: string, ...must: (string | RegExp)[]): Promise<void> {
const lines: Line[] = [];
await runCommand(cmd, (l) => lines.push(l), state);
const text = lines.map((l) => l.text).join("\n");
const missing = must.filter((m) => (typeof m === "string" ? !text.includes(m) : !m.test(text)));
const errored = lines.some((l) => l.tone === "err");
if (missing.length === 0 && !errored) {
passed++;
process.stdout.write(` ${GREEN}✓${RESET} ${cmd}${DIM} ${lines.length} lines${RESET}\n`);
} else {
const why = errored
? `emitted an error line: ${lines.find((l) => l.tone === "err")?.text}`
: `missing ${missing.map(String).join(", ")}`;
failures.push(`${cmd} — ${why}`);
process.stdout.write(` ${RED}✗${RESET} ${cmd}${DIM} ${why}${RESET}\n`);
}
}
process.stdout.write(`${BOLD}▸ browser shell${RESET}\n`);
await expect("help", "selftest", "gen", "tokenize", "quant", "mem", "gguf", "src");
await expect("selftest memory", "one buffer backs every view", /\d+\/\d+ checks passed/);
await expect("gen \"the cat\" -n 8", /8 tokens/, /tok\/s/, "stopped:");
await expect("gen \"j\" --json -n 30", "grammar");
await expect("tokenize the cat sat", "round-trip exact", "ids");
await expect("quant", "Q4_K", "144", "4.50");
await expect("quant Q4_K", "256 weights", "first 12:");
await expect("fp16 1.5", "bits", "out");
await expect("fp16 1e-8", "subnormal");
await expect("mem", "arena", "kv cache", "allocations");
await expect("model", "layers", "GQA", "vocab");
await expect("gguf", "magic ok", "tensors", "token_embd.weight");
await expect("src", "arena", "memory/arena.ts");
await expect("src arena", "stub for the headless run");
await expect("bench 8", /8 forward passes/, "ms/token", "MACs");
await expect("gpu", /WebGPU unavailable|adapter/);
// unknown commands must be reported, not thrown
{
const lines: Line[] = [];
await runCommand("nope", (l) => lines.push(l), state);
const ok = lines.some((l) => l.tone === "err" && l.text.includes("unknown command"));
if (ok) {
passed++;
process.stdout.write(` ${GREEN}✓${RESET} unknown command is reported, not thrown\n`);
} else {
failures.push("unknown command handling");
process.stdout.write(` ${RED}✗${RESET} unknown command handling\n`);
}
}
// every registered command must have help text
{
const bad = COMMANDS.filter((c) => !c.help || !c.name);
if (bad.length === 0) {
passed++;
process.stdout.write(` ${GREEN}✓${RESET} all ${COMMANDS.length} commands documented\n`);
} else {
failures.push(`undocumented: ${bad.map((c) => c.name).join(", ")}`);
process.stdout.write(` ${RED}✗${RESET} undocumented commands\n`);
}
}
const total = passed + failures.length;
process.stdout.write(`\n ${BOLD}${passed}/${total} shell checks passed${RESET}\n`);
if (failures.length) {
process.stdout.write(`\n${RED}${BOLD}failures:${RESET}\n`);
for (const f of failures) process.stdout.write(` ${RED}✗${RESET} ${f}\n`);
process.exit(1);
}The four-tab panel from earlier in this chapter is still worth a look for the memory view and the WebGPU benchmark:
Step 12
What's next
Three implementations of the same engine now: NumPy for clarity, C for the machine, TypeScript for the tab. The parts did not change. What changed each time was which costs the language makes visible. Python hides memory and punishes scalar loops, C hides nothing, JavaScript hides the collector until it stalls your UI.
That is the lesson of building it three times, and it is not really about inference. Pick the constraint you are optimising against, and the architecture follows.
Exercises
- 1Load a real model: wire a
<input type="file">toloadFromGGUFand point it at a small Q4_K GGUF. Everything needed is already above. Measure tokens/second and compare against the same model under S21's Python. - 2Add a
perfcommand toshell.tsthat reports per-layer timings, then find which of the seven matmuls dominates. Extendrun-shell-test.mtsto cover it: the rule in this chapter is that a command without a check does not exist. - 3Move the forward pass into a Web Worker so the main thread never runs a matmul. Note what has to change. The arena's ArrayBuffer becomes a
SharedArrayBuffer, which needs cross-origin isolation headers — a deployment constraint, not a code one. - 4Finish the GPU path: keep the weights quantised in device memory and dequantise inside the shader, then re-run
gpu. This is the fused/unfused argument from S08 again, one level down.
Check yourself
Questions
Answer before you look. A wrong answer tells you more than a right one, because the explanation points at the part of the chapter to reread.
Why does mini-picoLM allocate one ArrayBuffer up front and hand out views into it, rather than allocating tensors as needed?