It Runs End to End

01Rounding the real weights broke the model

stage 6 · the circuit was right, the model was not

The first thing I did was take the real weights of SmolLM2-135M, the model I had used so far, round them to ternary (−1 / 0 / +1), and stream them through the circuit. The circuit's answers matched numpy on all 239,232 words. The circuit works correctly.

As a model, though, it was completely broken.

187 times worse than picking at random. Changing the rounding per row barely helps. Each weight is off by 50–60%, and that error stacks up over 30 layers.

This stage turned up one more thing. The circuit ran at 100.0% of the datapath width, yet 83% of the time per token was fixed overhead from DMA calls. The format back then allowed "one x (input vector) per transfer," so every time the matrix changed I had to end the transfer, and one token took 181 calls. That became the homework for section 3.

02Switching to a model trained in ternary

stage 6.5 · the only one that fits in 512 MB

I looked for a model already trained in ternary that fits in this board's 512 MB. The only candidate was 1bitLLM/bitnet_b1_58-large (0.7B). I wrote the same quantization steps as the circuit in numpy and checked it: perplexity 17.90. Still ternary, and better than SmolLM2 in fp32 (18.31).

Its shape is different, though, so both the circuits and the way the data is packed had to be rebuilt.

To save memory, I dropped lm_head (the final matrix that scores all 32,002 vocabulary entries) to ternary. In INT8 the quality stays the same (17.90 → 17.89), but it doesn't fit in the region DMA can read (CMA). In ternary, perplexity doubles, 17.90 → 35.59, but the text doesn't fall apart. This is the quality price this setup pays. I come back to it in the last section.

03Fixing the handoff

stages 7 & 8 · a header, then a variable row length

How do you cut 181 calls? The circuit is already fast, so I changed how data is handed to it. Put a small header at the front of the stream saying "R rows follow," and the circuit can move on to the next group without ending the transfer when x changes. Next, I let the header also carry the length of one row.

Variable row length paid off in three ways at once. Less padding cut the bytes read by 17%. Group boundaries dropped to a quarter. And the cleanup step of adding partial sums on the CPU went away. One more thing: because the weights shrank to 141.3 MiB, attention could later share the same chip (section 5).

04Rebuilding attention

stages 7c & 9 · head_dim 96, no GQA, and a multiplier

The Stage 5b circuit only handles head_dim 64 and GQA 3:1, so it can't be used. Before rebuilding, I checked one more thing in numpy. Stage 5b squeezed the score into 8 bits with a fixed shift, s >> 8. That ignores the quantization scales of q and K. On the real model, the temperature of the exponent drifts from layer to layer and head to head.

I write m = TAU · sq · sk / √96, scaled by 2²⁴ into an integer, into the header. The circuit multiplies the score by it, then looks up the exponent table. Real values range 10,818–79,505. That doesn't fit in 16 bits, but fits exactly in the DSP's 18-bit input. That settles the previous entry's "the exponent table's temperature is a placeholder."

At 142.86 MHz, timing failed both times. The violation was not in my circuit but in the transfer length counter inside AXI DMA. I dropped the clock to 125 MHz. Two ports' datapath width is 2.000 GB/s, almost the same as the DDR wall of 2.009. The circuit is bandwidth-bound, so the lower clock costs almost nothing.

05Taking turns on one chip

stage 10 · four DMAs, no switch

I put the matmul circuit and the attention circuit in the same bitstream. Within a token, the two always run in alternation. So I built no switching logic. I placed four DMAs and just kick whichever one is needed. They never run at the same time, so they don't fight over DDR.

Sharing one bitstream cost 1.4% (projected 107.8 ms with separate bitstreams, 109.3 ms shared). Nearly all of that difference came from extra round trips, because DMA transfers were cut to 8 MiB or less. Resources: LUT 34%. Two-thirds of the chip is still free.

06Streaming a growing KV without moving it

stage 11 · scatter-gather instead of copying

To run end to end, K and V must grow by one row per token (the KV cache). This is where the attention circuit's input format became a problem.

The circuit takes a stream of [header][q][K, T rows][V, T rows], joined head after head. When T grows by 1, the start of V and the positions of every later head shift. A simple DMA can only read "one contiguous region," so the choices are to repack the whole KV every token (36 MB at T=512) or to end the transfer at every head (384 per token).

07The same answer on the Mac and the board

bit-identical logits on two machines

Running end to end means more to verify. Not just whether the circuit's answers are right, but whether the whole token is right, including the CPU's floating-point math. So I ran the same computation on the Mac too, and compared both sets of logits (the scores for all 32,002 vocabulary entries).

The Mac (Apple M4) and the board (Cortex-A9) are both ARM, so the same arm_neon.h instructions work on both. Even so, the results don't match unless you make them. These are the rules I set to make them match.

• Don't fuse multiply and add (-ffp-contract=off). Clang on the Mac fuses them into FMA by default
• Fix the order of sums (add in 4 lanes, then (0+1)+(2+3) at the end). No fast-math
• Write exp and rounding myself. expf and rintf can differ in the last bit from one libm to another
• Don't use vrecpsq_f32 to refine reciprocals. Same name, but on AArch64 it is a fused operation and gives different results from ARMv7

08Python was the slow part

the hardware was fine all along

When it first ran end to end, one token took 225.6 ms. That's twice the Stage 10 projection (109 ms). This table records knocking out the causes one at a time.

Stage 6 had the same story: "the circuit runs at 100% of the datapath width, yet call overhead is 83%." The same thing happened here. Handing data to the circuit takes as much work as making the circuit fast. That's where this stage spent most of its time.

09What grows as the context grows

only attention grows with T

I kept generating up to context 512 (the size of the board's KV buffer) and logged each component's time for every token.

Units are ms. Each row is the median of the surrounding 8 tokens. In a separate run under the same conditions, context 512 took 138.0 ms (7.25 tok/s).

10Projection vs. the real thing

where the budget went

At Stage 10 I measured the parts separately and added them up, projecting 109.2 ms at context 512. The real thing took 136.0 ms. Here is where the 26.8 ms difference came from.

11What I can and can't claim so far

what is measured and what is not

What I can claim. All of it measured on the real board.

• Given text, the board computes matmul and attention on my own circuits and generates a continuation
• Board and Mac logits match bit for bit across 191 tokens, and all 3,509 circuit calls match too
• At context 512, 136–138 ms / 7.3 tok/s (measured twice). 8.4 tok/s at short context
• Perplexity on the board through the circuits is 35.69. Float attention gives 35.56, so the circuit attention and fixed-scale KV cost 0.4%

What I can't claim.

• The biggest quality price is the ternary lm_head. The original fp16 gives perplexity 17.90; ternary doubles it to 35.59. It's a choice made to fit in CMA, not a circuit problem
• Generation is greedy (always the top-scoring word). Run it long and it repeats The family was very poor. That's common when you run a 0.7B model greedily, and numpy on the Mac produces the same text
• The CPU baseline is still only my own bench. I haven't compared against a practical implementation like llama.cpp
• The cause of the extra 250 µs (section 8) is unknown