← All notes and experiments

GPU concepts and data flow

Companion guide · Published 16 October 2026

Making a ternary reasoning model fast in vLLM

Published measurements are snapshots of the stated hardware, settings and revisions.

This guide explains the GPU ideas behind the series. Start with the basics, follow a small worked example, or look up a term while reading an experiment.

It is based on the standalone extract prepared on 15 October 2026, including the later same-day kernel-design follow-ups. The experiments have not been rerun for this article. The experiment reference preserves all 39 component-test records with their numerical checks and selected code.

Reading routes: Start with the basics, see the GPU data flow, or browse the experiments. For the complete-request measurements, use chapter 9.

Coming from the blog?

Find the concepts behind each chapter

The narrative chapters cover complete requests, power effects, speculative generation and task quality. This companion explains the component experiments. Their timings do not establish the complete-engine comparison in chapter 9, and they are not deployment-card measurements.

The basics

The model's inputs and calculations

A language model works with tokens: pieces of text represented by numbers, often smaller than a word. Inside the model, each token becomes a list of numbers—an activation vector.

A tensor is an array of numbers. A vector has one dimension; a matrix has two.

The model learns weights during training. A projection multiplies an activation vector by a weight matrix to make a new vector.

Each output is a dot product: multiply matching values, then add the products. The model repeats this work across its layers.

We'll use these letters throughout:

Symbol Meaning here
M Number of input token rows processed together
N Number of output values per token, and rows in the stored weight matrix
K Number of input values contributing to each output
X[M,K] Input activations
W[N,K] Stored weight matrix
Y[M,N] Output, calculated as X × Wᵀ

Prefill means processing the prompt's tokens. Decode means generating subsequent tokens, usually one new token per request at each step. A batch processes rows from several tokens or requests together. Speculative verification checks several proposed tokens together, which also creates a small batch of rows.

The GPU's workers and storage

A GPU divides work among streaming multiprocessors (SMs). A kernel is the function it runs; threads carry out the work.

On these NVIDIA GPUs, 32 threads form a warp. A thread’s position in it is its lane. One or more warps form a thread block and can cooperate. The scheduler chooses which ready warp runs next.

The data can live in several places:

Storage What to picture
Board memory, or DRAM The large store holding weights and other big arrays; also called device or global memory
L2 cache A smaller, faster store that can keep recently accessed data near the processing units
Registers Very small, fast storage for a running thread's values
Shared memory Limited on-chip storage that threads in a block can share

A load reads data into the computation; a store writes it back. Bandwidth is how many bytes can move per second. Latency is how long an operation takes, such as waiting for a particular load. You can have spare bandwidth and still be waiting for the next value you need.

The ordinary arithmetic units handle floating-point calculations, integer calculations and operations on bits. Tensor Cores are specialised hardware for small matrix products. DP4A is a different instruction: it multiplies four pairs of 8-bit integers and adds their products to a 32-bit integer sum. MMA means matrix multiply-accumulate: calculate matrix products and add them to a running answer. An accumulator holds that running answer; an operand is an input to an instruction.

Compressed numbers and their scales

A bit holds a 0 or 1; a byte contains eight bits. Quantisation represents a range of numbers using a smaller set of values. A scale maps those small stored values back to the size of the original numbers. Values sharing a scale form a quantisation group.

In this model's PQ2 weight format, each packed byte holds four two-bit codes. The decoded value is code−1, multiplied by the group's scale. A group of 128 weights has 32 code bytes and a 2-byte scale. Packed means several codes share a byte; unpacking or decoding extracts them. This use of “decode” is separate from the token-generation stage described above.

Q8 here means rounding activation groups into signed 8-bit integers with a scale. It's an approximation, so we check its numerical effect separately from its speed. INT8 and INT32 mean 8-bit and 32-bit integers. FP16, FP32 and FP64 are floating-point formats with 16, 32 and 64 bits. BF16, or bfloat16, is another 16-bit floating-point format with a different balance of range and precision. The code's dtype is the number format used by a tensor.

A tile is a small chunk of a matrix processed or stored together. A plane groups one kind of data, such as all code bytes or all scales. A mask selects particular bits or valid positions; a bit shift moves a code to the bit positions needed for an instruction. We'll draw these layouts in the terminology guide below.

Names for the model's other work

Attention compares the current token's query (Q) with stored keys (K), then uses the resulting weights to combine stored values (V). Together, the saved keys and values are the KV cache. Here, K in “KV” means keys; K in the matrix dimensions above means the input width. A head is one parallel attention calculation. Context length is the amount of token history available. Paged KV divides that history into storage chunks called pages.

QKV names the query, key and value projections, often joined into one projection. QK refers just to the query/key dot product.

Softmax turns attention scores into positive weights that sum to one. If attention is calculated in separate pieces, their partial results need the right softmax statistics to be merged correctly.

An MLP, also called a feed-forward network (FFN) here, is a set of projections and a nonlinear activation function. The gate/up projection expands and gates its inputs; down reduces them back to the model's working width. SiLU is the activation function used in the linked MLP tests. The vocabulary head, or LM head, projects the model's representation into scores for possible next tokens. This is different from an attention head.

Normalisation rescales values to control their magnitude. RMSNorm uses their root-mean-square magnitude. A residual carries an earlier representation forward so it can be added to a later result. The Hadamard transform mixes values through staged sums and differences; the butterfly example below shows how.

Gated DeltaNet (GDN) is the recurrent part of this hybrid model: it updates a stored state between steps rather than only reading an expanding attention history.

Reading the experiment names and results

Use this glossary when an experiment introduces an unfamiliar term. You can also go straight to the worked examples.

Programming languages and matrix kernels

GEMV means matrix–vector multiplication; GEMM means matrix–matrix multiplication. The MMQ name labels the quantised matrix-multiply implementations used in these experiments. CUDA is NVIDIA's GPU programming platform; Triton is the programming language/compiler used for many of these kernels. A compiler translates source code into instructions the GPU can run.

Specialisation means compiling a version tailored to known inputs such as a matrix size; a runtime value is supplied when the program runs rather than fixed during compilation. PyTorch creates tensors and calls the GPU operations. ctypes lets Python call compiled functions directly.

Attention implementations and kernel selection

FA2 means FlashAttention 2, an attention implementation. FlashInfer provides another set of attention implementations. A backend is the chosen implementation of an operation. Dispatch is the rule that chooses a kernel for the current shape. A shape is the set of array dimensions.

A layout says how those logical positions are arranged in memory; the same shape can have several layouts. Fusion combines operations into one kernel so they can share intermediate work or data. Native, reference, baseline and control mean the comparison implementation in that particular test; they don't all refer to one fixed version across the report.

Saved inputs, timing and temporary storage

A capture saves inputs or outputs for later checks. A standalone replay runs those saved tensors through the component again. A CUDA graph records GPU operations so they can be launched together with less repeated CPU work. A warm-up runs the operation before timing so compilation and initial setup don't dominate.

First-call timings include costs that warmed measurements may leave out. Scratch or workspace means temporary storage. A buffer is an allocated region of memory; allocation reserves that storage. A stream orders GPU operations. A finite-output check rejects infinities and NaNs, the special values used for invalid or undefined numerical results.

Comparisons and validation checks

In paired timings, we compare a candidate with a control under similar conditions, often alternating their order. A screen tries several candidates; confirmation gives promising candidates a closer check. A gate is a stated pass/fail rule. Poisoning outputs means filling a result buffer with invalid values before a replay, then checking that the operation really overwrites them.

A mutation check changes an input to make sure the replay uses its new value. A linked replay feeds one component's result into the next, such as gate/up → SiLU → down. It still doesn't include every operation in the model.

Unless stated otherwise, speedup = control time ÷ candidate time: 2× means half the time; 0.8× means slower. Matching output bits or a small numerical error is a different check from task accuracy.

Speedup and numerical error

Unless stated otherwise, speedup = control time ÷ candidate time: 2× means half the time, and 0.8× means slower. Percentage gains use the definition in the relevant test. They can't be added across experiments with different baselines. Bit-exact means the stored output bits agree.

Relative RMS error measures the size of a difference relative to the reference: the square root of the summed squared differences divided by the summed squared reference values. RMSE is the square root of the mean squared difference, without that reference-size normalisation.

Neither is a task-accuracy score.

Profiling tools and GPU instructions

A profiler measures what the GPU does while it runs; Nsight Compute is the profiler used for the hardware counters here. A counter records a hardware event or activity measure. SASS is the GPU's machine-instruction listing. PRMT rearranges bytes in registers.

NVTX labels regions for profiling. cp.async copies data asynchronously into shared memory. A barrier makes cooperating threads wait until the required participants reach a point. A deterministic reduction combines partial answers in a fixed order so thread arrival order doesn't change the result.

Units

Units: ms means milliseconds, µs microseconds, GB/s billions of bytes per second, and GiB multiples of 2³⁰ bytes. In context-length labels such as 4K, K is shorthand for roughly four thousand tokens, not the matrix's input-width symbol.

Tiles, layouts and moving data

If tiles and memory layouts are new to you, start here. The examples use tiny arrays so you can see what moves where. Real tests use much bigger tiles. A compute tile, a storage tile, and a quantisation group serve different purposes even when they cover some of the same values.

Term Plain-English meaning Relevant tests
Tile and tile geometry A chunk of work handled together 5, 11, 13, 15, 16, 22
Row-major and tile-major Different orders for placing the same values in memory 1, 4, 24, 25
Stride, alignment and coalescing How addresses line up with efficient GPU loads 1, 2, 4, 24
Warps, lanes and reductions Which threads cooperate, and how they combine answers 11, 15, 16, 17
Butterfly and shuffle Staged exchanges between pairs of values or threads 1, 12, 17, 18
Planes, fragments and Tensor Cores How packed values are arranged for matrix instructions 6, 7, 21, 23
Occupancy, spills and pipelining Keeping hardware busy without exhausting local storage 14, 19, 20, 23

Warp, lane, block and reduction

A thread executes instructions and holds register values. A warp is a group of 32 threads; a lane is one thread's position within that warp. A thread block contains one or more warps and can share on-chip shared memory. A multiprocessor, or SM, runs resident blocks and schedules their warps.

A reduction combines many values into fewer values, often by addition. A projection's dot product is a multiply followed by a sum. Different threads can calculate pieces, then reduce their partial sums. Floating-point addition is rounded, so changing the summation order can change the last bits even when the formula is mathematically identical.

Split-K divides the input/reduction dimension among multiple blocks. This creates more parallel work when there are too few output tiles. The blocks then need a way to combine their partial answers; the tested deterministic variants use a separate fixed-order reduction. Attention splitting has a similar parallelism goal but needs a softmax-aware merge, not just addition of independently normalised outputs.

Tile: a small rectangle of a large calculation

A tile is a chunk of a matrix that a kernel processes together. Instead of asking each thread to calculate one entire answer independently, a group of threads cooperates on a tile and reuses the input data it loaded.

For these projections, X[M,K] holds input tokens, W[N,K] holds weights, and the result is Y[M,N] = X × Wᵀ. A tile of the result might cover two tokens and four output values per token:

               Output columns N
             0  1  2  3 | 4  5  6  7
Token 0      [*  *  *  *]| .  .  .  .
Token 1      [*  *  *  *]| .  .  .  .
             -----------+-----------
Token 2       .  .  .  . | .  .  .  .
Token 3       .  .  .  . | .  .  .  .

[*] = one 2-token × 4-output tile
Each answer still sums over K inputs.

The kernel usually works through K a chunk at a time, adding each chunk's contribution to the answer. Tile geometry means these chunk sizes and the thread arrangement. In the code, BM or BT often describes token rows, BN describes weight/output rows, and BK describes input elements. BC in the integer GEMV tests counts 128-element scale groups, so BC=4 covers 512 input elements per loop step.

Always read the particular function's arguments: these names are conventions, not universal hardware units.

Larger tiles can reuse a weight across more tokens, but must keep more partial answers alive. Smaller tiles create more independent blocks but may reload the same inputs. A tail is the leftover edge when dimensions do not divide evenly; masking prevents padded lanes from reading or writing invalid values. A kernel can process tiles without changing the matrix's stored layout.

Row-major versus tile-major: which values are neighbours in memory?

Row-major stores one complete row, then the next. Tile-major stores one rectangular tile, then the next tile. The numbers and logical matrix stay the same; their physical addresses change.

For a 4×4 matrix divided into 2×2 storage tiles:

Logical matrix       2×2 tile boundaries
 A B C D              A B | C D
 E F G H              E F | G H
 I J K L             -----+-----
 M N O P              I J | K L
                      M N | O P

Memory, read left to right:

Row-major:
 A B C D  E F G H  I J K L  M N O P

Tile-major, row-major inside each tile:
 A B E F  C D G H  I J M N  K L O P

Tile-major is useful when a warp needs pieces of several rows at the same time: those pieces can be placed together instead of far apart. It is not a transpose—the logical row and column of each value remain unchanged. The kernel's address calculation must match the chosen storage order.

The actual layout in tests 24–25 is more specialised than this toy example: it groups code bytes by output-row tile, quantisation group and lane, so neighbouring lanes can load adjacent 16-byte pieces. The layout within a tile is part of the design; there is no single universal tile-major format.

Prepared planes are a separate change. Moving all code bytes into one plane and all scales into another does not by itself make weights tile-major. Tests 1 and 4 separate codes from scales; tests 24 and 25 investigate how code bytes are ordered for a warp.

Stride, contiguous memory, alignment and coalescing

Stride is the address jump required to move one position along an array dimension. If adjacent values occupy adjacent storage, that dimension is contiguous. Taking every second row creates a larger row stride even if each individual row remains contiguous. Reshaping or permuting dimensions can change the view of data; a subsequent .contiguous() can copy it into a new physical order, as the tile-major conversion does.

Alignment means an address starts at a suitable byte boundary for a load. The raw PQ2 format repeats a 2-byte scale plus 32 code bytes—a 34-byte block. That stride is inconvenient for uniformly aligned wide loads. Separate planes can make the code stream easier to align.

Coalescing concerns a whole warp's addresses, not just one thread's load width:

Four example lanes    Addresses requested

Nearby pieces:
 lane 0               bytes   0–15
 lane 1               bytes  16–31
 lane 2               bytes  32–47
 lane 3               bytes  48–63

Separated pieces:
 lane 0               bytes    0–15
 lane 1               bytes 1024–1039
 lane 2               bytes 2048–2063
 lane 3               bytes 3072–3087

The first pattern can use memory transactions more efficiently. The real GPU warp has 32 lanes; the number of transactions also depends on alignment, request size and architecture. A wide load in each lane is not enough if the lanes still request scattered regions. Conversely, row-major storage can be efficient when the kernel requests suitable contiguous chunks.

Butterfly: mix pairs, then mix more distant pairs

A butterfly is a staged pattern of pairwise exchanges. In the Hadamard transform used here, each pair (a,b) becomes (a+b,a−b). The next stage pairs values farther apart. With four inputs, the pattern is:

Start:       [ a, b, c, d ]

Stage 1: pair neighbours (distance 1)
             [ a+b, a−b, c+d, c−d ]
                u0   u1   u2   u3

Stage 2: pair across halves (distance 2)
             [ u0+u2, u1+u3,
               u0−u2, u1−u3 ]

Normalized Hadamard: divide by sqrt(4).

For a power-of-two length B, this takes log2(B) stages: distances 1, 2, 4, and so on. In the Hadamard code, i ^ h selects the partner by flipping one bit of the index. This spreads each input's contribution across the outputs using additions and subtractions, rather than a general dense matrix multiply. The model also uses sign factors around this transform.

A shuffle is an instruction that lets one lane obtain a register value from another lane in the same warp. A butterfly shuffle, such as the SASS instruction family SHFL.BFLY, uses an XOR-style partner pattern. Shuffles can implement butterflies, rearrange operands or help reductions; they are not themselves the mathematical Hadamard transform.

In test 1, counts of butterfly-shuffle instructions measure lane-exchange work in the packed-dot implementation. More such instructions do not mean the model performed more Hadamard transforms. In test 18, the Hadamard transform is an actual preparation step being fused. A downward-shuffle sum, as used by some CUDA reductions in test 17, is another exchange pattern and should not be confused with the full sum-and-difference transform.

Quantisation groups, code planes and register fragments

A quantisation group shares a numerical scale. In these PQ2 weights, 128 values use 32 bytes of two-bit codes plus one 2-byte scale. This is a numerical/storage boundary, not necessarily a compute tile: a kernel tile can cover several such groups.

A packed byte contains four two-bit codes:

Bit positions:      7 6 | 5 4 | 3 2 | 1 0
Packed codes:        c3 |  c2 |  c1 |  c0

Plane 0: take c0 from each packed byte
Plane 1: take c1 from each packed byte
Plane 2: take c2 from each packed byte
Plane 3: take c3 from each packed byte

Plane order arranges corresponding input values so each extracted code plane can multiply the matching inputs efficiently. These code planes are different from the separate code-versus-scale planes described above. The raw code represents a weight value of code−1; the masked-plane kernels account for this offset using an activation-sum correction.

RT, or row tiling in the later CUDA code, is the number of 16-row weight tiles one warp handles together.

A fragment is the portion of a Tensor Core matrix tile assigned to one thread's registers. The hardware matrix instruction expects those values in a specific arrangement across lanes. Test 23 loads matching fragments directly into registers to avoid the shared-memory staging used by the earlier Triton route. MMA means multiply-accumulate: multiply matrices and add their products to an accumulator. DP4A is a smaller integer dot instruction, not a Tensor Core matrix instruction.

Occupancy, spilling, pipelining and persistent kernels

Occupancy measures resident warps relative to the SM's supported maximum. Here, resident means the warp has space on the SM; it may still be waiting for an input. Registers and shared memory are finite: a block that uses more leaves room for fewer other blocks. Resident doesn’t mean ready to execute; a warp can still be waiting for a load or a dependency. Higher occupancy can hide waits, but is not automatically faster.

A register cap limits how many registers the compiler may use per thread. A spill moves values that cannot remain in registers to device-backed local memory, which can be cached. This adds load/store work and can erase a gain from an otherwise attractive tile. Shared memory is explicitly shared within a block; a register spill is not simply a transfer into shared memory.

Pipelining overlaps stages, such as fetching the next tile while calculating the current one. Prefetch requests data early. Unrolling exposes multiple loop iterations to the compiler at once. All can create more independent work, but also keep more values alive. Tests 19–20 investigate this trade-off.

A persistent kernel launches a pool of blocks that loop over multiple work tiles instead of launching a separate block for every tile. It changes work assignment and scheduling; it doesn’t mean the weights stay permanently in registers. Test 17 explores whether this improves how the last tiles fill the GPU.

Two profiling phrases that crop up later

A long-scoreboard stall means a warp is waiting for an outstanding memory-related result it needs before it can continue. No eligible warp means the scheduler has no ready warp to issue an instruction. Neither measure is a direct count of wasted seconds; their denominators and collection method matter.

A wave is a round of blocks that can fit on the GPU at once. If only a few blocks remain for the final wave, some SMs have no work. A persistent grid uses a fixed pool of blocks to work through several tiles, which can change that balance.

Staging means putting values in an intermediate storage location before they're used. A pipeline stage is one step or live buffer in an overlapped load/compute sequence. An asynchronous staging ring cycles through buffers so one can be filled while another is consumed.

What good GPU data flow looks like—and where our model gets awkward

A GPU works best when nearby threads read nearby data, reuse it and have more work ready while they wait. The arithmetic is fast; supplying and arranging its inputs is often the harder job.

For our Bonsai 27B PQ2 workload on RTX 3090/3090 Ti GPUs, we want to keep weights compressed in memory while handing them to the hardware in a useful order, often for just one or a few tokens. This overview explains the design pressures behind the standalone tests; it does not introduce new measurements.

What a well-fed GPU would do

A useful way to picture the ideal is a short cycle:

  1. Load a tile with a few aligned memory requests.
  2. Decode values into the order the instructions need.
  3. Reuse them across several tokens and keep partial answers close to the compute units.
  4. Write the finished output once.

Meanwhile, load the next tile and give other multiprocessors independent work.

Process diagram: Weights in GPU memory; Nearby, aligned loads; Decode into the required operand layout; Reuse tile across tokens; Compute and accumulate in registers; Write final outputs; Load next tile early

That's the aim. In practice, some values need shared memory or a trip back to board memory. Shared memory is useful when it enables cooperation or reuse; it becomes overhead when it is only a staging stop that could have been avoided. Tensor Cores are useful when the matrix shape and operand arrangement let them perform enough useful work. Other instruction paths can be better for very small batches.

There are three separate ways to fall short:

Limit What it means in plain English What might help
Memory bandwidth Too many bytes must pass through the memory system Compression, reuse, fewer intermediate writes, efficient addresses
Memory latency or dependencies Threads wait for particular inputs even though aggregate bandwidth has headroom Independent work, earlier loads, suitable tile sizes and more ready warps
Compute or instruction throughput Arithmetic, unpacking, shuffling or address calculations occupy execution capacity Better instructions, less rearrangement, appropriate Tensor Core tiles

These limits can coexist. High memory utilisation doesn’t mean every warp receives its next operand immediately, and high GPU activity doesn’t mean every instruction contributes equally to the final answer.

Why prompt processing and token generation want different kernels

During prompt processing (prefill), many input-token vectors are available together. A weight tile can contribute to many token outputs before it is discarded. This creates a substantial matrix multiply, which can make good use of Tensor Cores.

During single-request generation (decode), the model usually has one new token vector to process at each step. It still needs the large projection matrices, but now each weight contributes to far fewer token outputs. For an ordinary autoregressive request, the next token's input depends on the preceding result, so a kernel can’t simply fill its token tile with future tokens.

Multiple independent requests or speculative verification can supply extra rows, but their sizes differ from long prefill batches.

Same weight tile; different amount of useful reuse

Prefill:
 load weights once → token 0, token 1, ... many tokens

One-request decode:
 load weights once → one token
 next token needs the preceding result

Small batch / verification:
 load weights once → a few token rows
 may need padding to fit a hardware tile

This is why the report contains separate floating GEMV, DP4A, prefill MMQ and small-batch Tensor Core experiments. The ideal dispatch chooses a suitable kernel for the available rows; one universally large tile would waste work on small batches.

Where our model makes this harder

1. The weights are small in memory but need work before arithmetic. Each PQ2 group stores 128 weights as 32 code bytes plus a 2-byte scale: 34 bytes, or 2.125 bits per weight including the scale. The GPU matrix instructions used here cannot directly consume that stored PQ2 block as a dense matrix tile.

The kernels extract codes, interpret code−1, arrange integer or floating operands, and apply scales. Compression saves memory capacity and traffic, but moves decoding and arrangement work into the kernel. Expanding the entire model to dense weights would give back much of that saving.

See storage and decoding tests and compute-path tests.

2. The full weight set cannot remain in the fast on-chip caches. The real-weight replay includes 6.33 GiB of distinct weights, far beyond these GPUs' L2 capacity. A benchmark that repeatedly uses one cached matrix can hide the cost of bringing different layers' weights from board memory.

So the useful target is to stream the needed weights efficiently and reuse each loaded tile while it is available, rather than assume the whole model stays close to the cores. See test 13.

3. Compact storage does not guarantee efficient loads. The original 34-byte blocks interleave scales and codes. Even after separating those planes, neighbouring threads can still request small pieces from widely separated rows. The ideal layout puts the bytes a warp consumes together next to each other.

The load-only probe in test 24 measures about 417–447 GB/s for the current gate/up pattern versus 951 GB/s for a tile-major pattern. Those are code-load probe figures, not complete projection throughput. Larger contiguous row-major reads also improve the probe: tile-major is one solution, not the only one.

4. There may be too little independent work even with a good layout. A projection with fewer output rows produces fewer row tiles. Making each tile larger can reduce the block count further, leaving some multiprocessors without work. Splitting the K reduction creates extra parallel blocks, but their partial answers must then be combined.

In test 25, the first tile-major MMA variant helps gate/up while down regresses without split-K. Efficient addresses and sufficient parallelism are separate requirements. Tests 30–31 later add split-K and recover the down-projection gain.

5. Reuse and prefetch compete for limited working storage. Keeping more tokens' partial sums, more weight groups or the next tile in flight consumes registers and sometimes shared memory. Beyond a point, fewer warps fit on an SM, or temporary values spill into device-backed local memory.

The ideal keeps enough independent work ready to cover delays without carrying an excessive live working set. Tests 19–20 show why deeper pipelines and stricter register caps cannot be assumed to help.

6. Attention has a second large input: the token history. The full-attention component uses 24 query heads, 4 KV heads and head dimension 256. As context grows, it reads more keys and values, stored in pages with the actual model's strides.

One query may need its history divided among many blocks to expose enough parallelism. Those blocks must merge softmax statistics correctly. Compressing KV reduces its payload but adds scales, conversion and numerical choices. The ideal serves history reads efficiently while keeping split/merge overhead proportionate to context length.

See attention tests.

7. A decode step has several different jobs. Projections are interleaved with normalisation, Hadamard transforms, MLP activation and, in this hybrid model, attention or recurrent Gated DeltaNet state updates. Outputs become later operations' inputs, creating dependencies and intermediate tensors. Recurrent state also has to advance consistently between steps.

Fusion can remove some memory round trips, but only where shapes, dependencies and numerical behaviour permit it. The standalone projection replays explicitly omit these state updates; they cannot establish complete-step performance. See fusion test 18 and the scope notes on each replay.

8. A faster data path still has to calculate an acceptable answer. Matching a different reduction order can change floating-point rounding; changing Q8 group size changes the approximation itself. Even a tiny difference before quantisation can change an integer code.

A useful optimisation keeps the agreed numerical behaviour, and the test must say whether it checks exact bits, a projection-error threshold, or something broader. This report's local checks do not establish task-level quality. See quantisation controls.

How the experiments fit the ideal

Ideal property Experiments that investigate it
Stream bytes in the order the warp needs them Prepared layouts, access-pattern probes
Reuse each weight for more useful outputs Batched arithmetic, small-batch MMA
Put operands directly where instructions expect them Code planes and register fragments
Have enough independent tiles and ready warps Geometry, stall diagnosis and pipelining
Avoid unnecessary intermediate reads and writes Fused preparation
Scale parallel history processing to context length Attention splitting and KV formats
Keep numerical changes explicit and tested Group-size controls and each test's validation

That gives us a way to read the tests: follow the data, see where it waits, and ask what each change costs in storage or extra work. The same change can help one shape and hurt another.