Hardware · Transformer Inference

An LLM Transformer Decoder in SystemVerilog

A synthesisable pre-norm decoder block — Q8.8 fixed-point, 4×4 systolic MAC array, KV-cache, 84-test verification suite, Vivado synthesis to Artix-7. Walk through every module. Compare to H100, Trillium, Cerebras, Groq, and Sohu. Trace the architecture back to "Attention Is All You Need".

SystemVerilog Pre-norm decoder Q8.8 fixed-point Systolic MAC KV-cache RSqrt LUT + Newton–Raphson PWL softmax Vivado / Artix-7 CocoTB · iverilog
Repo → Contents Report Companion: MMUL deep-dive
01

What this deck covers

Twenty-six slides. Half walk through every RTL module in the repo at the level of a code review; the other half put the design in context — the algorithm, the historical hardware, and the real shipping silicon that runs the same workload on a million-times bigger budget.

  1. Why a transformer-decoder accelerator exists
  2. The decoder block, end to end
  3. Q8.8 fixed-point — the only number format in the design
  4. Module: transformer_pkg.sv — the whole arithmetic library in one file
  5. Module: processing_element.sv — the MAC tile
  6. Module: systolic_array.sv — 4×4 array with fill / drain accounting
  7. Module: softmax_unit.sv — PWL exp + reciprocal-by-LUT FSM
  8. Module: layer_norm.sv — division-free norm via RSqrt + Newton–Raphson
  9. Module: feed_forward.sv — two GEMMs and a ReLU
  10. Module: multi_head_attention.sv — heads, masking, softmax, KV-cache
  11. Module: transformer_decoder.sv — the top-level FSM
  12. Streaming variant — 99.4% fewer registers
  13. KV-cache: kv_cache_bram.sv and friends
  14. Verification — three tiers, 84 tests
  15. Synthesis on Artix-7 — what's tight, what's slack
  16. A short history of transformer hardware
  17. Real systems — NVIDIA Hopper Transformer Engine
  18. Real systems — Blackwell, Trillium, Trainium, MI300
  19. Real systems — Cerebras, Groq, Tenstorrent, Etched, SambaNova, Apple ANE
  20. The KV-cache problem at scale — PagedAttention, NVFP4 KV
  21. Q8.8 vs INT8 vs FP8 vs FP4 — when fixed-point wins
  22. Transformer-on-FPGA — FTRANS, EdgeBERT, FQ-BERT, ViTA
  23. Where this design fits
  24. Interactive — decoder cycle counter
  25. Interactive — Q8.8 fixed-point explorer
  26. Limitations & what would change at the next scale
02

Why a transformer-decoder accelerator exists

"Attention Is All You Need" (Vaswani et al., NeurIPS 2017, arXiv:1706.03762, posted 12 June 2017) replaced the recurrence and convolution stack with self-attention plus position-wise MLPs. It exposed a compute pattern that has dominated silicon roadmaps ever since:

By 2020, attention had become a measurable fraction of inference time, and the academic-hardware community responded with attention-as-a-primitive papers (A³ HPCA 2020, SpAtten HPCA 2021, ELSA ISCA 2021, FACT ISCA 2023). Industry caught up in 2022 when NVIDIA shipped the Hopper Transformer Engine — the first time a GPU had per-layer FP8 statistics built in. Blackwell Ultra (2025) doubled the throughput of the SFU's MUFU.EX2 instruction, the one softmax actually sits on — silicon admitting that softmax is finally a first-class concern.

What this repo is

This is the educational analogue of all of the above, scaled to fit in 16 GB of host memory and on a $200 Artix-7 FPGA. Every block — GEMM, softmax, LayerNorm, KV-cache, top-level FSM — is small enough to read in one screen of code. None of the moves above are skipped: the softmax is a real PWL approximation, the LayerNorm is a real division-free RSqrt path, the KV-cache really avoids re-projection of past tokens.

Scope

D_MODEL = 64, N_HEADS = 4, D_FF = 256, MAX_SEQ_LEN = 128, Q8.8 16-bit signed fixed-point, 4×4 PE systolic array. Vivado 2025.2 synthesises the streaming variant on Artix-7 (xc7a35tcpg236-1). The point isn't the throughput — it's that all eight RTL files together are smaller than a single CUTLASS kernel.

03

The decoder block, end to end

One transformer_decoder instance corresponds to one decoder layer. Stacking N of them gives an N-layer decoder; weight memories can be shared across layers (area-optimal, time-multiplexed) or replicated (throughput-optimal). The dataflow per token is:

  token_emb ──► [LayerNorm 1] ──► [Multi-Head Attention] ──►(+)──► [LayerNorm 2] ──► [Feed-Forward] ──►(+)──► out_emb
                                                              ▲ Residual 1                                ▲ Residual 2
                  └────────────────────────────────────────────┘             └─────────────────────────────┘

This is a pre-norm decoder: LayerNorm sits before each sub-layer, with a residual that bypasses the sub-layer. Pre-norm is the architecture used by GPT-2, GPT-3, GPT-4, LLaMA-1/2/3, Mistral, Qwen, and most modern decoder-only LLMs. The original "Attention Is All You Need" Transformer was post-norm; the switch happened around 2019 because pre-norm trains far more stably at large scale (Xiong et al., "On Layer Normalization in the Transformer Architecture", ICML 2020).

Per-token sequence of operations

  1. LayerNorm 1 on the input embedding — zero mean, unit variance, learnable γ/β.
  2. Multi-head attention: compute Q from current token; pull K, V for all prior positions from the KV-cache; compute scores Q·Kᵀ / √d_k; apply causal mask; softmax; weighted sum over V; concat heads; project with WO.
  3. Residual 1: token_emb + attention_out (saturating add in Q8.8).
  4. LayerNorm 2 on the residual.
  5. Feed-Forward: ReLU(x·W₁ + b₁) · W₂ + b₂. The "wide-then-narrow" 4× expansion that all transformer FFNs use.
  6. Residual 2: residual1 + ffn_out.
  7. Update KV-cache for this position; advance seq_pos.

Module hierarchy

transformer_decoder          ← Top-level FSM (12 states)
├── layer_norm × 2           ← Pre-attention & pre-FFN normalisation
├── multi_head_attention     ← Causal MHA with KV-cache + time-multiplexed softmax
│   └── softmax_unit         ← Shared, MAX_SEQ_LEN-wide
├── feed_forward             ← Two-layer FFN with ReLU
└── transformer_pkg          ← Parameters · types · FP utilities
04

Q8.8 fixed-point — the only number format in the design

Every value in the datapath — activations, weights, biases, partial sums, softmax outputs, normalised activations — is 16-bit two's-complement Q8.8. That means 8 integer bits, 8 fractional bits, signed. Range −128.000 to +127.996, resolution 1/256 ≈ 0.0039.

PropertyValue
FormatQ8.8 signed (1 sign · 7 integer · 8 fraction); width 16
Range[−128.0, +127.996]
Resolution2−8 = 0.00390625
Multiplier16 × 16 → 32-bit signed; one Artix-7 DSP48E1 slice
Accumulator32-bit, full precision; truncated to 16-bit Q8.8 at module boundaries
Saturating addfp_sat_add in transformer_pkg.sv — clamps to ±0x7FFF / 0x8000
Multiplyfp_mul: (a·b) >>> FRAC_BITS — arithmetic right-shift of the 32-bit product

Why fixed-point at all?

Where Q8.8 hurts

Q8.8 has the same 65 536-level dynamic range as INT16 — about 1/40 of FP16's range. For D_MODEL ≥ 1024 and unscaled activations (the GPT-2-large era and beyond), per-tensor Q8.8 starts to fail because individual attention activations hit the rails. That's the regime the industry left fixed-point behind for: SmoothQuant (per-channel), AWQ (per-group), GPTQ, FP8-with-per-tile-scale (the Hopper Transformer Engine), and now NVFP4 (per-block FP4 with E4M3 scale).

This design lives in the small-D regime where Q8.8 works, on purpose. The package-level types make it the single point of change: edit DATA_WIDTH, FRAC_BITS, and the FP utility functions, rebuild, re-run the 84-test suite.

05

Module: transformer_pkg.sv

233 lines. The whole arithmetic library in one place — every other module imports it. Edit this file to scale the design.

Parameters

ParameterDefaultComment
D_MODEL64Embedding dimension
N_HEADS4Heads (D_HEAD = D_MODEL/N_HEADS = 16)
D_FF256FFN inner dimension (4× expansion)
MAX_SEQ_LEN128Context window
VOCAB_SIZE512Used in the bigger top-level only
DATA_WIDTH / FRAC_BITS16 / 8Q8.8
ACC_WIDTH32Full-precision MAC accumulator
PE_ROWS / PE_COLS4 / 4Systolic array shape

The four utility functions every other module uses

rtl/transformer_pkg.sv — fp_mul, fp_sat_add
// Multiply two Q8.8 values, return Q8.8 (truncated)
function automatic signed [15:0] fp_mul(input signed [15:0] a, input signed [15:0] b);
  logic signed [2*DATA_WIDTH-1:0] product;
  product = a * b;
  return to_data(product >>> FRAC_BITS);
endfunction

// Saturating add for Q8.8
function automatic signed [15:0] fp_sat_add(input signed [15:0] a, input signed [15:0] b);
  logic signed [DATA_WIDTH:0] sum;
  sum = {a[DATA_WIDTH-1], a} + {b[DATA_WIDTH-1], b};
  if (sum > $signed({1'b0, {(DATA_WIDTH-1){1'b1}}}))      return {1'b0, {(DATA_WIDTH-1){1'b1}}};
  else if (sum < $signed({1'b1, {(DATA_WIDTH-1){1'b0}}})) return {1'b1, {(DATA_WIDTH-1){1'b0}}};
  else return to_data(sum[DATA_WIDTH-1:0]);
endfunction

The interesting one: fp_inv_sqrt

A division-free reciprocal-square-root, ~12 bits of accuracy in Q8.8. The same shape of code (CLZ → normalise → 32-entry LUT → one Newton–Raphson step → denormalise) appears three times in the design: here for LayerNorm's 1/√σ², in softmax_unit for 1/Σexp, and as a pattern that students can transplant to any Q-format reciprocal they need. The key trick is CLZ-normalise once, denormalise at the end — it lets a 32-entry LUT cover the full 16-bit input range to identical accuracy across decades.

For odd leading-zero counts the code multiplies by a Q2.14 constant for √2 (= 23170) so the binary-only shift can resume after the LUT step. This is verbatim what x86's RSQRTSS and ARM's VRSQRTE.F do internally — minus the IEEE rounding modes.

06

Module: processing_element.sv

52 lines. The smallest possible MAC tile — three flops, one multiplier, one adder. Same shape as Kung & Leiserson 1978; same shape as a single TPU MXU lane in 2025. The pattern is forty-eight years old and still right.

rtl/processing_element.sv — the whole module
always @(posedge clk or negedge rst_n) begin
  if (!rst_n) begin
    a_out <= 0; w_out <= 0; accumulator <= 0;
  end else if (clear) begin
    a_out <= 0; w_out <= 0; accumulator <= 0;
  end else if (enable) begin
    a_out       <= a_in;                            // skew east
    w_out       <= w_in;                            // skew south
    accumulator <= accumulator + (to_acc(a_in) * to_acc(w_in));
  end
end

assign acc_out = accumulator;

What's interesting

Synthesis result on Artix-7 (xc7a35t, 100 MHz target)

MetricValue
LUTs34
FFs64
DSP48E11
Fmax229.9 MHz

One DSP, one PE. The 4×4 array uses 16 DSPs — under 30% of the xc7a35t's 90 DSP48E1 slices. Headroom for an 8×8 array is real on this part.

07

Module: systolic_array.sv

92 lines. Generates ROWS × COLS PEs, wires their A-east and B-south chains, and tracks a completion counter so the host can tell when the last accumulator has settled.

Topology

A[0,*]→ A[1,*]→ A[2,*]→ A[3,*]→ B[*,0]↓ B[*,1]↓ B[*,2]↓ B[*,3]↓ acc[r,c] PE 0,0PE 0,1PE 0,2PE 0,3 PE 1,0PE 1,1PE 1,2PE 1,3 PE 2,0PE 2,1PE 2,2PE 2,3 PE 3,0PE 3,1PE 3,2PE 3,3

A streams left→right (blue), B streams top→bottom (purple), accumulators stay put (green).

Latency accounting

Why output-stationary, not weight-stationary?

This design loads a fresh weight tile for every projection (Q, K, V, O, FFN1, FFN2), so weights aren't reused enough to amortise weight-stationary's load cost. Output-stationary is also the natural shape for the QKᵀ score computation, which dominates per-head cycles. NVDLA, Apple ANE, and most edge NPUs go weight-stationary only because they hold a whole layer's weights in SRAM for batched inference — a regime this small reference design doesn't aim at.

08

Module: softmax_unit.sv

332 lines — the largest module in the repo by line count. Implements the standard max-subtract-exp-normalise recipe, every step of which is approximated:

  S_IDLE ─► S_FIND_MAX ─► S_SUBTRACT_EXP ─► S_SUM ─► S_RECIP ─► S_NORMALIZE ─► S_DONE

Step 1 — find max (numerical stability)

One pass over the input vector, holding the running maximum. Standard textbook trick: subtracting max(x) before exp prevents overflow without changing the result, because softmax is shift-invariant.

Step 2 — PWL exp approximation

A 4-segment piecewise-linear approximation of exp(x) over the negative range — that's all softmax ever sees after subtracting max:

rtl/softmax_unit.sv — approx_exp
function automatic signed [15:0] approx_exp(input signed [15:0] x);
  if (x >= 0)               return 16'sh0100;                      // exp(0) = 1.0
  else if (x > -16'sh0100)  return 16'sh0100 + x;                  // (-1, 0): linear
  else if (x > -16'sh0200) begin                                 // (-2, -1)
    result = 16'sh0060 + (x + 16'sh0100) >>> 1;                  // ~0.375 + 0.5*(x+1)
    return (result > 0) ? result : 16'sh0004;
  end else if (x > -16'sh0400) begin                          // (-4, -2)
    result = 16'sh0020 + (x + 16'sh0200) >>> 2;                  // ~0.125 + 0.25*(x+2)
    return (result > 0) ? result : 16'sh0004;
  end else                  return 16'sh0004;                      // floor for very negative
endfunction

This is good enough for token-selection ordering (which is all that matters for greedy/argmax decoding) and produces probabilities that sum to ≈ 1.0 within a few LSBs in Q8.8. The TB confirms Σprobs ≈ 0xFE / 0x100 for an 8-element uniform input.

Step 3 — sum the exps, then reciprocal

Computing p_i = exp(x_i) / Σexp would need a divider. Instead the unit computes recip = 65536 / exp_sum in one shared cycle, then multiplies each exp_i by recip with the same fp_mul the rest of the design uses. Because fp_mul implicitly does (a·b) >> 8, the formula works out to (exp_i << 8) / exp_sum — exactly the original division.

The reciprocal itself uses the same CLZ-normalise → 32-entry Q2.14 LUT → one Newton–Raphson step → denormalise pattern as fp_inv_sqrt. Newton–Raphson here is r ← r·(2 − x·r), which converges quadratically; ~12 bits of accuracy after one step is enough for Q8.8 outputs.

Synthesis result

MetricValueComment
LUTs537Mostly the 32-entry case-LUT and the Newton step
FFs584FSM, exp_vals[], shifted[], idx
DSP48E13Reciprocal mul + Newton mul + per-element normalise
Fmax47.9 MHzCritical path: combinational compute_reciprocal

The 47.9 MHz figure is one of the two reasons the full design doesn't meet the 100 MHz target. Pipelining the reciprocal across two FSM states would more-than-double Fmax — the deferred work is documented in README.md.

09

Module: layer_norm.sv

151 lines. Computes yi = γi · (xi − μ) / √(σ² + ε) + βi across a VEC_LEN-element vector. The interesting part is what's not here: any divider.

  S_IDLE ─► S_MEAN ─► S_VARIANCE ─► S_NORMALIZE ─► S_DONE

Three division-replacement tricks in one module

The +1 epsilon (literal 16'sh0001 = 1/256 in Q8.8) is enough to keep the LUT index away from zero on the constant-input case (variance = 0). The behavioural test Constant→zero exercises this: feed a vector of identical values, expect zero output.

Why this matters

A hardware integer divider for 32-bit values is ~600 LUTs and 12+ pipeline stages on Artix-7. Replacing it with a 32-entry LUT + two 16×16 multiplies (Newton–Raphson) keeps the path entirely inside the DSP fabric. The cost is ~12 bits of accuracy — more than enough for Q8.8.

Synthesis result

MetricValueComment
LUTs1,585VEC_LEN = 64 is the dominant cost — centered[] array
FFs2,157Per-element registers across the FSM
DSP48E17γ·(x−μ)·invstd is two cascaded muls per element
Fmax29.2 MHzCritical path: fp_inv_sqrt Newton–Raphson chain

29.2 MHz is the slowest module in the design and therefore the system-level limiter. The path is the combinational chain inside the Newton–Raphson step: 16-bit multiply → subtract from constant → 16-bit multiply → variable shift. All inside a single FSM state. Pipelining across two FSM states is the documented path to ≥ 100 MHz.

RMSNorm vs LayerNorm

Modern LLMs (LLaMA, Qwen, Mistral) use RMSNorm instead of LayerNorm — same RSqrt path, but no mean-subtraction. RMSNorm is roughly 30% cheaper and slightly faster to train. Migrating this module to RMSNorm is mostly deletion: drop S_MEAN, drop the centered[] array, keep fp_inv_sqrt.

10

Module: feed_forward.sv

125 lines. Two linear layers and a ReLU. The wide-then-narrow shape every transformer FFN uses — D_MODEL=64 expands to D_FF=256, then collapses back to D_MODEL=64.

  S_IDLE ─► S_LINEAR1 ─► S_RELU ─► S_LINEAR2 ─► S_DONE
              (D_FF iters)      (parallel)   (D_MODEL iters)

Why FFN is the biggest GEMM in the layer

For one decoder layer at sequence length S, the FLOP budget breaks down (per "Attention Is All You Need" §3.3 and the standard expansion):

StageFLOPsFraction
Q, K, V projections (3×)6 · S · D²~17%
Output projection WO2 · S · D²~6%
QKᵀ scores2 · S² · D≪ for small S
scores · V2 · S² · D≪ for small S
FFN expand (4D × D)8 · S · D²~22%
FFN contract (D × 4D)8 · S · D²~22%

FFN is roughly 44% of the layer for typical S, and at single-token inference (S = 1) it's the dominant compute. This is why GPT-3, GPT-4, and LLaMA all keep the 4× FFN ratio — and why FlashAttention, which optimises the attention block, has so little to say about FFN.

The ReLU choice (not SwiGLU)

Modern LLMs use SwiGLU (Shazeer 2020): three weight matrices, two for the gate, one for the projection, sigmoid(W₁x) ⊙ (W₂x) → W₃ projection. SwiGLU outperforms ReLU on training perplexity but costs 3 GEMMs and a sigmoid per token instead of 2 GEMMs. The reference design uses plain ReLU because:

What's missing for a "real" FFN

11

Module: multi_head_attention.sv

362 lines. The single biggest module in the repo. Implements the canonical scaled-dot-product attention from "Attention Is All You Need" §3.2 — minus the divide-by-√d_k, which is folded into pre-scaling of the weights.

  Attention(Q,K,V) = softmax(Q·Kᵀ / √d_k) · V    over each of N_HEADS in parallel

Ten-state FSM

StateCyclesWhat happens
S_IDLE1Wait for start
S_PROJ_QKV3 · D_MODELCompute fresh Q, K, V vectors for current token
S_WRITE_CACHE1Append K, V at seq_pos in the cache
S_SCOREN_HEADS · seq_pos · D_HEADCompute Q · Kᵀ scores over all past positions
S_SOFTMAX_PREP1Pad scores beyond seq_pos with −128.0 (causal mask)
S_SOFTMAX_RUN~5·MAX_SEQ_LENRun shared softmax_unit for this head
S_SOFTMAX_STORE1Capture probabilities
S_WEIGHTED_SUMN_HEADS · D_HEAD · seq_posCompute Σ pi · Vi per head
S_OUTPUT_PROJD_MODEL²Concat heads, project with WO
S_DONE1Pulse valid

The shared softmax unit

One softmax_unit instance is time-multiplexed across heads. The instance is sized for VEC_LEN = MAX_SEQ_LEN, and positions beyond seq_pos are padded with the smallest representable Q8.8 value (−128.0) so the PWL exp approximation forces them to ~0. This implements the causal mask as a number, not as a control bit — saves a separate masking pipeline.

The KV-cache interface

K and V for all prior positions are held in dual-port memory. On every attention step:

In the streaming variant the cache is kv_cache_bram — a real dual-port BRAM that does element-by-element vector writes and single-element reads. In the non-streaming variant it's a flat unpacked array, intended only for simulation.

What's not here

12

Module: transformer_decoder.sv

249 lines. The conductor — instantiates LayerNorm-1, MHA, LayerNorm-2, FFN, threads them with two residual additions, and runs a 12-state top-level FSM.

  S_IDLE ─► S_LN1_START ─► S_LN1_WAIT ─►
            S_ATTN_START ─► S_ATTN_WAIT ─► S_RESIDUAL1 ─►
            S_LN2_START ─► S_LN2_WAIT ─►
            S_FFN_START ─► S_FFN_WAIT ─► S_RESIDUAL2 ─► S_DONE

Things the conductor handles that the sub-modules don't

Per-token cycle budget

Approximate, for D_MODEL=64, N_HEADS=4, seq_pos = 64 (mid-context):

StageCycles
LayerNorm 1~200
MHA Q/K/V proj~200
MHA score + softmax + weighted sum~2,500
MHA output proj~4,100
LayerNorm 2~200
FFN linear 1 + ReLU + linear 2~16,400
Total per token~24,000 cycles

At the achievable 29.2 MHz that's ~1.2 ms per token, ~830 tok/s — for a one-layer 64-dim toy model on a $200 FPGA. Stack 12 layers at this clock and you get ~70 tok/s, a reasonable rate for a small embedded model. The streaming variant trades throughput for area and is what synthesises within the 16 GB Vivado memory budget.

13

Streaming variant — 99.4% fewer registers

The repo ships two top-levels:

Register-bridge (high-throughput)

  • transformer_decoder_top.sv
  • Pipeline registers at sub-module boundaries
  • One token, one pass; full-width registers everywhere
  • Designed for back-to-back token throughput
  • Reference; doesn't fit on a small Artix part as-is

Streaming (minimum-area)

  • transformer_decoder_top_stream.sv
  • Single shared datapath, time-multiplexed
  • BRAM-backed weight memories (weight_bram, kv_cache_bram)
  • −99.4% register count vs the bridge variant
  • The intended Artix synthesis target

Why two?

For an educational repo, having both makes the trade-off concrete: the streaming variant is the bridge variant with the boundary registers stripped and the weights moved to BRAM. Diff one against the other and you can see, line by line, what "minimum-area" means in practice. The streaming variant's regression test (tb_transformer_decoder_stream.sv) produces bit-identical outputs to the bridge variant for the same weight load.

From stream_results.log:

--- Test 1: Single token at position 0 ---
  out_emb[0] = 0097 (0.5898)   ← matches bridge variant exactly
  out_emb[1] = ff7f (-0.5039)
  ...
[PASS] KV cache active: token-1 output differs from token-0
============================================
  Results: 4 PASSED, 0 FAILED
============================================

The streaming pattern, in one sentence

Replace combinational weight arrays with single-port BRAM and an address generator; replace pipeline registers between sub-modules with a single bus + handshake; let the FSM serialise everything that used to overlap. Trade ~10× fewer registers for ~3–5× more cycles per token. On Artix-7 that's the difference between "synthesises" and "exceeds 16 GB host memory".

14

KV-cache: kv_cache_bram.sv

Without a KV-cache, autoregressive generation re-projects every prior token's K and V on every step — O(n²·d) per token. With a KV-cache, only the fresh Q is computed; past K, V are read from memory — O(n·d) per token. The streaming variant uses a real dual-port BRAM (bram_dp) sized for MAX_SEQ_LEN × D_MODEL.

Address mapping

One BRAM each for K and V. Linear address: mem[pos · D_MODEL + dim] — row-major, position-major.

Two-port pattern

PortUsePattern
A (write)Append K or V at current seq_posPulse wr_start, the module auto-walks D_MODEL elements from the input vector
B (read)Per-element fetch by (pos, dim)One-cycle latency, used by the score / weighted-sum loops

Sizing

For MAX_SEQ_LEN = 128, D_MODEL = 64, Q8.8: 128 × 64 × 16 bits × 2 (K and V) = 32 768 bits = 4 KB. Fits in two RAMB18E1 BRAMs comfortably (the xc7a35t has 50 of them).

By contrast, a real LLM:

ModelKV cache @ FP16KV cache @ NVFP4 (Blackwell)
LLaMA-2-7B, 4K context~2 GB~0.25 GB
LLaMA-2-70B, 4K context~40 GB (~size of weights)~5 GB
LLaMA-3.1-405B, 128K context~640 GB~80 GB

That's the entire reason long-context LLM inference is HBM-bandwidth-bound rather than compute-bound. PagedAttention (vLLM, SOSP 2023) borrows OS virtual memory to cut KV waste from 60–80% to under 4%. NVIDIA has been pushing FP8 and NVFP4 KV quantisation to claw the size back. Same idea here, fifteen orders of magnitude smaller.

15

Verification — three tiers, 84 tests

Tier 1 — Behavioural

54 tests, ~2 s on a laptop. scripts/verify_behavioral.py implements a bit-accurate Python model of the RTL: same Q8.8 representation, same truncation rules, same RSqrt LUT entries. No simulator required. The fastest path to confidence in algorithmic correctness.

Tier 2 — RTL simulation

30 tests in tb/sv/, run by scripts/run_sim.sh under iverilog 12.0. PE, systolic array, softmax, BRAM, full decoder, streaming decoder. The integration tests reproduce known-good outputs from a Python golden model.

Tier 3 — CocoTB

Python-driven RTL co-simulation in tb/cocotb/. Currently exercises PE and softmax — the modules where Python's expressiveness pays for itself in stimulus generation and golden-model comparison.

What the behavioural tier covers

ModuleTestsKey checks
Fixed-point utilities16Q8.8 roundtrip · multiply with ±1 LSB tolerance · saturating add at the rails
Processing element9Single MAC · accumulate · forwarding · clear · 20-op random vs Python golden
Systolic array82×2 matmul producing [[19,22],[43,50]] exactly · single-element · clear
Softmax unit8Uniform · dominant · ordering · sum ≈ 1 · back-to-back
Layer norm5Constant→zero · symmetry · γ/β scaling · centering
Feed-forward4ReLU zeroing · bias propagation
Decoder integration4Full pipeline · KV-cache · 2-token sequential

Recent fixes — git log as a verification record

That's six bugs caught after the initial release. Each one became a TB; each TB is in the regression suite.

Structural lint — scripts/lint_check.py

Walks all eight RTL files, checks module/endmodule balance, package imports, reset patterns (posedge clk or negedge rst_n), and cross-file instantiation resolution. Cheap to run, catches typos that iverilog's parser elides.

16

Synthesis on Artix-7 — what's tight, what's slack

Vivado 2025.2, target xc7a35tcpg236-1, 100 MHz clock constraint. The full top-level exceeded the host's 16 GB memory budget, so sub-modules were synthesised individually. The numbers below are per-sub-module and aren't additive (the array shares logic with the PEs it contains).

ModuleLUTsFFsBRAMDSP48E1Fmax (MHz)Comment
processing_element346401229.9Headroom for 2× clock
systolic_array (4×4)455901016156.8Skew-flop chain limits Fmax
softmax_unit5375840347.9Critical: combinational reciprocal
layer_norm1,5852,1570729.2Critical: Newton–Raphson chain
transformer_decoder_top (estimated)~2,611~3,7060~27≤ 29.2System-level Fmax bound by layer_norm

Why 100 MHz wasn't met

Two modules contain combinational logic inside a single FSM state that's too long for a 10 ns period:

This isn't a bug; it's a deliberate trade. The combinational implementation is one screen of code; the pipelined version is two — and an educational design is allowed to leave the pipeline as homework. The README documents the gap.

What it costs on a 7-series xc7a35t

Total budget: 20 800 LUTs, 41 600 FFs, 50 BRAM, 90 DSP48E1.

17

A short history of transformer hardware

YearEvent
1978Kung & Leiserson, Systolic Arrays for VLSI. The PE/array shape that this repo's processing_element/systolic_array modules still implement.
2015Google TPU v1 — 256×256 INT8 systolic, deployed for inference SLA.
2017Vaswani et al., "Attention Is All You Need" (NeurIPS, arXiv:1706.03762). Transformer is born.
2017Microsoft Project Brainwave — FPGA "soft NPU" on Stratix 10, 39.5 TFLOPS at sub-ms latency in MS-FP8.
2017NVIDIA V100 ships first Tensor Core (4×4×4, FP16 mul / FP32 acc) — pre-transformer in target, but the GEMM accelerator that won the transformer era.
2018Habana Goya HL-1000 — first commercial deep-learning inference accelerator. CNN-and-BERT-class.
2018Google Edge TPU — 4 TOPS INT8 at 2 W. Pure CNN; no useful transformer support.
2020A³ (HPCA) — Ham et al., the first major paper to architect attention as a hardware primitive. ELSA (ISCA 2021), SpAtten (HPCA 2021), FACT (ISCA 2023) follow.
2020NVIDIA A100 — TF32, BF16, FP64 TC, 2:4 sparsity, 312 TFLOPS BF16. The first GPU on which 100B-parameter models trained in production.
2020FTRANS (Li et al., ISLPED) — block-circulant BERT on FPGA, 8.8× energy improvement vs GPU.
2021EdgeBERT (Tambe et al., MICRO) — entropy-based early exit + adaptive attention span on a 12 nm ASIC.
2021FQ-BERT (Liu et al., DATE) — fully integer/fixed-point BERT including softmax and LayerNorm in the integer domain. Direct ancestor of this repo's algorithmic moves.
2022FlashAttention (Dao et al., NeurIPS, arXiv:2205.14135) — reduces attention HBM traffic to O(N²d²/M); 2–4× speedup with no approximation. Hardware-software co-design lesson of the decade.
2022NVIDIA H100 + Transformer Engine — per-layer FP8 statistics, 1979 TFLOPS FP8 dense.
2024FlashAttention-3 (Shah et al., arXiv:2407.08608) — Hopper-aware, 740 TFLOPS FP16, ~75% peak utilisation, ~1.2 PFLOPS FP8.
2024NVIDIA Blackwell B200 — FP4, NVFP4 (16-element block, FP8 E4M3 scale); 20 PFLOPS FP4.
2024Google TPU v6e Trillium — MXU enlarged from 128×128 to 256×256, first resize since 2017.
2024Etched Sohu announced — first transformer-only ASIC.
2025NVIDIA Blackwell Ultra doubles SFU MUFU.EX2 throughput. The first GPU generation where softmax is treated as a first-class architectural concern.

The throughline: every generation, the silicon learns one more thing about transformers. Pre-2020 it learned "you need a lot of MACs". 2022 it learned "you need 8-bit numbers and per-layer scales". 2024 it learned "you need an explicit memory primitive (TMA, TMEM)". 2025 it learned "softmax is on the critical path; the SFU has to be wider". Forty-eight years from Kung's PE to the Blackwell Ultra SFU — the primitive doesn't change, the surroundings do.

18

Real systems — NVIDIA Hopper Transformer Engine

Hopper (2022) was the first GPU architecture marketed by reference to a non-GEMM workload. The Transformer Engine isn't a separate IP block — it's a software runtime + hardware features that turn each forward/backward pass into a per-layer scaled FP8 matmul.

What the Transformer Engine actually does

  1. Profile the activation/weight tensor's max-abs in FP16.
  2. Pick a per-layer scaling factor so the tensor fits FP8's range. Two FP8 formats:
    • E4M3: 4-bit exponent, 3-bit mantissa, range ~ ±448 — used for forward weights and activations, where precision matters.
    • E5M2: 5-bit exponent, 2-bit mantissa, range ~ ±57 344 — used for gradients, where range matters.
  3. Cast the tensor to FP8, multiply on the 4th-gen Tensor Core, accumulate in FP32.
  4. Cast back to BF16/FP16 for downstream non-GEMM ops (LayerNorm, softmax, residual).

The two enabling primitives

How the reference design compares (very loosely)

FeatureH100 Transformer EngineThis repo
DatatypeFP8 E4M3 / E5M2 + BF16 accQ8.8 mul + 32-bit acc
Per-layer scaleYes, runtime-managedNo — fixed-point only
GEMM topology4×8×16 reduction-tree TC × 132 SMs4×4 output-stationary systolic × 1
Async DMATMANone
SoftmaxSFU MUFU.EX2 + warp shuffle reductionFSM, PWL exp, reciprocal LUT
NormVector lane, FP16/BF16 divFSM, RSqrt LUT + Newton–Raphson, no divider
Peak FP8 dense1 979 TFLOPS~3 GMAC/s @ 29 MHz

The ratio is roughly 10⁹. The architectural moves are the same — block GEMM, max-subtract softmax, normalise, residual. Implementing them in 1 500 lines of SystemVerilog instead of 50 million transistors is the entire point.

Common confusion

"H100 does 989 TFLOPS" cites BF16. The FP8 figure is 1 979 TFLOPS dense or 3 958 sparse. Tech press articles routinely compare H100 BF16 to AMD/TPU FP8 — apples vs oranges.

19

Real systems — Blackwell, Trillium, Trainium, MI300

NVIDIA Blackwell B200 (2024)

  • 4NP TSMC, dual-die, 192 GB HBM3e, 8 TB/s
  • 5th-gen Tensor Core, 20 PFLOPS FP4 dense
  • NVFP4: 16-element block + FP8 E4M3 scale + per-tensor FP32 scale
  • TMEM — 256 KB/SM tensor-private memory
  • Blackwell Ultra (2025): 2× SFU EX2 throughput for softmax

Google TPU v6e Trillium (2024)

  • 4 nm, MXU enlarged 128×128 → 256×256
  • BF16 / INT8, ~926 TFLOPS BF16
  • SparseCore for embedding lookups (LM head)
  • OCS pod fabric — runtime-reshapeable
  • v7 Ironwood (2025) adds FP8 → ~4.6 PFLOPS

AWS Trainium2 (2024)

  • NeuronCore-v2: 4 heterogeneous engines
  • Tensor (matmul) · Vector · Scalar · GpSimd
  • ~667 TFLOPS BF16 / FP8
  • 96 GB HBM3e, NeuronLink fabric

AMD MI300X / MI355X

  • 192 GB HBM3 / 256 GB HBM3e
  • MI300X: 1.3 PFLOPS FP16, 2.6 PFLOPS FP8
  • MI355X: FP4 / MXFP4 support, ~5 PFLOPS FP4
  • Same per-block scale story as NVIDIA's NVFP4

Three different bets

Notice what no shipping system has: a hardwired fixed-point datapath of the kind in this repo. Production silicon picks per-layer scaled FP8/FP4 instead because (a) range adapts, and (b) the area cost of FP arithmetic is no longer dominant at 4 nm. Q8.8 lives on in educational and edge-FPGA niches — this design is squarely in the first.

20

Real systems — the rest of the field

DeviceTopologyDatatypePeak / SRAMTransformer-specific feature
Cerebras WSE-3 (CS-3)900 K cores, full 300 mm wafer, 5 nmFP16, BF16, FP8125 PF · 44 GB SRAM, 21 PB/sWeight Streaming + MemoryX (4 TB–2.4 PB) + SwarmX gradient reduction
Groq LPU (TSP)Functional slices, 14 nm, 900 MHzFP16, INT8188 TF FP16 / 750 TOPS INT8 · 230 MB SRAM, no DRAMCompiler-scheduled deterministic everything — autoregressive decoding has constant-shape steps; static schedule wins
Tenstorrent Wormhole / Blackhole80 / 140 Tensix cores, 12 nm / 6 nmFP8/16, BFP2/4/8292 / 745 TF FP8 · 210 MB SRAM (Blackhole)Per-core matrix + vector + RISC-V, NoC mesh — dataflow that maps each transformer layer to N cores
Tesla Dojo D1354 cores/chip × 25 chips/tile, 7 nmFP32, BF16, CFloat89 PF / training tile · 11 GB SRAM/tileCFloat8 — user picks the exp/mantissa split per layer (a per-tensor analogue of Hopper's per-layer FP8 scales)
Etched Sohu (announced 2024)Hardwired transformer dataflow, 4 nmFP8 (vendor claim)vendor: 500 K tok/s Llama-70B / 8-Sohu server · 144 GB HBM3eWhole transformer block hardwired — no flexibility for non-transformer architectures
SambaNova SN40LRDU, 1040 PCU + 1040 PMU CGRABF16 / FP8638 TFLOPS · 520 MiB SRAM + 64 GiB HBM + 1.5 TiB DDRStreaming dataflow — compiler fuses hundreds of operators across one pipeline
Apple Neural Engine (M4)16-core, weight-stationaryFP1638 TOPS · shared SoC SRAMNative scaled-dot-product-attention layer; ml-ane-transformers reference reports 10× speedup, 14× lower peak memory
Graphcore IPU (Mk2 / Bow)1472 tiles, 14 nm / wafer-on-wafer 7 nmFP16, FP32250 TFLOPS FP16 · 900 MB SRAMHardware stochastic rounding at full clock — no FP32 master-weight copy needed

Three philosophies, one workload

The reference design in this repo lives at the educational end of the third axis: it does only a transformer decoder and nothing else, and the hardwiring is its strongest feature for teaching purposes.

21

The KV-cache problem at scale

The KV-cache is the single biggest difference between transformer training and transformer inference. Training is GEMM-bound; inference, especially at long context, is HBM-bandwidth-bound — and the bandwidth is consumed by KV reads.

The arithmetic

Per token, per layer, the decoder reads 2 · seq · D_MODEL bytes of KV and computes 2 · seq · D_MODEL MACs. The arithmetic intensity is O(1) FLOPs per byte — the canonical memory-bound case. No matmul accelerator helps. On an H100 at single-batch inference, MMUL utilisation drops to 1–3%; the rest of the chip is idle waiting on HBM.

The two industry fixes

PagedAttention (vLLM, SOSP 2023)

Borrow OS virtual memory: split the KV cache into fixed blocks, give each request a block table mapping logical to physical. Cuts wasted KV memory from 60–80% to under 4%, gives 2–4× throughput over FasterTransformer/Orca at the same latency. Hardware implication: attention kernels must do gather-scatter from a block table.

KV quantisation

FP8 (E4M3 / E5M2) KV halves memory at near-FP16 accuracy. NVFP4 KV halves it again — <1% accuracy loss on MMLU-PRO, MBPP, LiveCodeBench, Ruler 64K (NVIDIA blog, 2025). LMDeploy ships per-head per-token asymmetric INT4/INT8.

What this design does instead

For MAX_SEQ_LEN = 128, D_MODEL = 64, Q8.8: the entire KV cache is 4 KB. It fits in two BRAMs. There is no HBM, no PagedAttention, no quantisation pressure — the problem doesn't exist at this scale. What survives is the interface: append-on-token, broadcast-on-attention. That interface is the same on this Artix-7 design and on a 1 000 000 000 000-parameter model running on 16 H100s.

22

Q8.8 vs INT8 vs FP8 vs FP4 — when fixed-point wins

FormatBitsRangeLSB / GranularityMul costRange/Precision split
Q8.8 (this repo)16±1282⁻⁸ ≈ 4·10⁻³1 DSP48Static — fixed point
INT8 + per-tensor scale8 + 16-bit scalescale-dependentscale·2⁻⁷~½ DSPStatic per tensor
INT8 + per-channel scale8 + 16-bit/chper-channelper-channel~½ DSP + scale mulPer-output-channel — used by SmoothQuant, AWQ
FP8 E4M3 (Hopper fwd)8±4482⁻⁹ near 1.0~⅓ DSP eq.Per-element float
FP8 E5M2 (Hopper grad)8±57 3442⁻⁵ near 1.0~⅓ DSP eq.Bigger range, less precision
NVFP4 / MXFP4 (Blackwell)4 + shared E4M3 scale per 16±6 per element~½ near 1.0tiny mul + shared scale fan-outBlock-scaled — best of both

What this matters for

The area / energy curve

Multiplier area is roughly quadratic in mantissa width. So:

The curve eventually meets a wall: at FP4, multiplier energy is so low (~0.03 pJ at 4 nm) that operand movement dominates the budget. That's the regime block scaling exists for — share one E4M3 scale across 16 elements, broadcast it once, save 15× on scale-fetch energy.

23

Transformer-on-FPGA — the genealogy

YearProjectWhat it didHeadline number
2017MS Project BrainwaveFPGA "soft NPU" on Stratix 10, MS-FP839.5 TFLOPS, sub-ms latency
2020FTRANS (Li et al., ISLPED)Block-circulant weight compression for BERT/RoBERTa16× model-size reduction · 27× perf vs CPU · 8.8× energy improvement vs GPU
2021EdgeBERT (Tambe et al., MICRO)Entropy-based early exit, sentence-granularity DVFS, FP quantisation; 12 nm ASIC53× energy savings vs Jetson TX2
2021FQ-BERT (Liu et al., DATE)Fully integer/fixed-point BERT — softmax and LayerNorm in the integer domain, all intermediates fixed-point28.91× perf-per-watt vs CPU/GPU, < 1% accuracy loss on SST-2
2022ViTA (Nag, Datta et al.)Vision-transformer accelerator on Zynq ZC7020~90% utilisation @ 0.88 W, 150 MHz
2022FlashAttention (algorithm)Tile-and-fuse Q·Kᵀ + softmax + scores·V into a single SRAM-sized kernel2–4× speedup, no approximation
2023–2024ALSAP, DOTA, FAMOUS, FNM-TransSparse attention, N:M sparsity, length-adaptive pipelining; UltraScale+ targets—

The fixed-point lineage

This repo's core algorithmic moves all trace back to the FQ-BERT line: integer softmax (PWL exp), integer LayerNorm (RSqrt LUT + Newton–Raphson), integer attention scoring, integer FFN. The FQ-BERT paper showed that you can run BERT entirely in the integer domain with < 1% accuracy loss — this design implements the same idea at smaller scale and with simpler approximations.

Why FPGAs are still relevant in the 2026 landscape

24

Where this design fits

PropertyThis repoFTRANS (2020)FQ-BERT (2021)ViTA (2022)H100 (2022)
TargetArtix-7 xc7a35tVCU118VCU118Zynq ZC70204 nm GPU
Clock29.2–229.9 MHz (mod-dep.)~200 MHz~200 MHz150 MHz1.78 GHz
NumericQ8.8 fixed-pointBlock-circulant 16-bitMixed integer/FXPINT8FP8 + per-layer scale
D_MODEL64768 (BERT-base)768768any (compiler tiles)
Compute4×4 systolic + FSM softmax/LNMultiple SAs + dedicated softmax/LNReconfigurable, INT softmax/LNSingle dense engine132 SMs × Tensor Cores + SFUs
KV-cacheDual-port BRAM, append-on-token———HBM3, PagedAttention via vLLM
VerificationPython golden + iverilog + CocoTB, 84 testsPer-paperPer-paperPer-paperInternal QA + CUDA test stack
Lines of RTL~1 500largelargelarge(unpublished)

Honest scope statement

This is not a competitive transformer accelerator. It will not run LLaMA-7B; the systolic array is two orders of magnitude too small, the KV-cache is six orders of magnitude too small, the clock is one order of magnitude too slow. It is an educational reference implementation with three stated goals:

  1. Make every architectural primitive of a transformer decoder small enough to read in one screen of code.
  2. Show, by construction, that softmax + LayerNorm + KV-cache + GEMM compose into one autoregressive step.
  3. Provide a regression-tested baseline for changing one variable at a time — RMSNorm, GQA, FlashAttention-style fusion, INT8 with per-channel scales, an 8×8 array — without rewriting the others.

If you're reading the RTL because you're learning hardware design, that's the audience. If you're reading it because you want to deploy a transformer to an FPGA, look at FQ-BERT or ViTA instead and bring the patterns from this repo to your own scale.

25

Interactive — decoder cycle counter

Choose a configuration; see the per-token cycle budget for one decoder layer at the chosen sequence position. Numbers come from the FSM cycle counts in the actual modules.

64
4
4
64
29
LayerNorm
—
QKV proj
—
Score+Softmax+WS
—
FFN
—
Total cycles / token
—
Latency (ms)
—
Throughput (tok/s)
—
FFN as % of total
—

Notice how FFN's share of the budget climbs as D_MODEL grows (because FFN is O(D²) while the rest is O(D · seq)) and how seq_pos has almost zero effect on per-token cost at small D — the KV-cache is doing exactly the work it's there to do.

26

Interactive — Q8.8 fixed-point explorer

Type a value; see how Q8.8 stores it, the rounding error, and the worst-case error across the format's range. Compare with FP8 E4M3 (Hopper forward), FP8 E5M2 (Hopper backward), and FP4 E2M1 (Blackwell).

Q8.8 (this repo) INT8 BF16 FP8 E4M3 FP8 E5M2 FP4 E2M1
Stored bits
—
Decoded value
—
Abs error
—
Range
—

Try 0.5898 in Q8.8 — that's the actual out_emb[0] from the integration test (0x0097). Try 1e-4 in FP4 (rounds to zero) versus Q8.8 (also zero — Q8.8 is no better at small numbers than FP4 is). Try 200 in Q8.8 (saturates) versus FP8 E4M3 (representable). The format choices are full of these quiet trade-offs.

27

Limitations & what would change at the next scale

What this design doesn't do

What the next scale-up would touch

The point of all of this

None of those changes are mysterious. They're all in the literature, in shipping silicon, or in the companion AI_MMUL_Unit repo. The reference design is small enough to make each change a self-contained patch, with the regression suite catching anything that drifts. That's the role of educational RTL — not to replace H100, but to make H100's architectural choices legible.

Closing

From Kung's 1978 PE to the Blackwell Ultra SFU, the primitive doesn't change — the surroundings do. This repo is the smallest place where you can see the whole loop close: a 16-bit MAC, a piecewise-linear exp, a Newton–Raphson reciprocal square root, a KV-cache, and an FSM that schedules them so a token comes out the other end.