Kernels: Where the Time Actually Goes

A model is a graph of operators, and at runtime that graph is a sequence of kernels, two or three of which account for eighty per cent of the time. At some point you stop reasoning about the graph and start reasoning about a loop nest. That is the job I spend most of my time on at Arm: taking one operator, working out what fraction of the machine it actually uses, and finding the restructuring that closes the gap.

This chapter is that process from first principles: the most obvious matrix multiply, its memory traffic, and the roofline bound from Part 1; then the rebuild; then the same for convolution, which has four genuinely different algorithms and no universal winner. Along the way we hit the most useful practical lesson I know: cutting FLOPs and cutting latency are different projects.

One illustrative core is used throughout, consistent with Part 2. It is not any particular product, and every number below follows from these parameters.

Everything is single-core unless stated.

The Naive Triple Loop and Its Memory Bill

Matrix multiply, the way everyone writes it first. \(A\) is \(M \times K\), \(B\) is \(K \times N\), both row-major.

for (int i = 0; i < M; i++)
  for (int j = 0; j < N; j++)
    for (int k = 0; k < K; k++)
      C[i*N + j] += A[i*K + k] * B[k*N + j];

Correct, three lines, and at \(M = N = K = 1024\) it runs at about two GFLOP/s on our core, six per cent of peak. The slowness is predictable from the access pattern before you run anything.

Counting the traffic

The work is \(F = 2MNK = 2.147 \times 10^9\) FLOP. The compulsory traffic, the bytes you would move with perfect caches, is the three matrices touched once: \(3 \times 1024^2 \times 4 = 12.58\) MB. That is 171 FLOP/byte against a ridge point of 2.5, so matrix multiply is compute bound by a factor of about seventy and should run at peak.

Now count what the loop moves. A[i*K + k] walks contiguously: one row is 4 KB, sits in L1, and is reused across every j, so it is negligible. B[k*N + j] walks down a column with a stride of \(N \times 4 = 4096\) bytes, so every access lands on a different 64-byte line: 1024 lines per inner loop, 64 KB of line traffic to deliver 4 KB of useful floats. Sixteen floats share a line, so the next sixteen values of j would reuse them, except that 64 KB is exactly the size of L1 and a row of \(A\) is streaming through it too. The working set thrashes, so assume each column walk re-fetches from L2:

That is an effective intensity of 0.031 FLOP/byte. We took an operation with an intensity of 171 and implemented it 5500 times worse, landing far to the left of the ridge.

What the roofline allowed

If L2 delivers 32 bytes per cycle, 64 GB/s at 2.0 GHz, then 68.7 GB takes 1.07 s against a compute-bound 67 ms.32 bytes per cycle of sustained L2 read bandwidth is reasonable for a mid-range core, but varies by a factor of two across designs. Measure yours with a strided read microbenchmark sized to fit L2 and miss L1. The naive loop is sixteen times slower than the machine allows, achieving \(2.147 \times 10^9 / 1.07 = 2.0\) GFLOP/s, 6.3 per cent of peak, which is what you measure in practice.

Some of that sixteen is free: reorder the loops to i, k, j, both B and C become unit-stride, and B traffic falls to about 4 GB.Interchange alone typically takes this kernel from two GFLOP/s to eight or ten. Most compilers will not do it at -O3, because proving it legal needs alias analysis that a raw pointer signature defeats. Adding `restrict` sometimes helps; writing the loop in the right order always does. Then we do the clever things.

Takeaway: Matrix multiply has a compulsory intensity of about 170 FLOP/byte and should be strongly compute bound. The naive loop implements it at an effective 0.03 because the column walk of \(B\) touches a new cache line per element. That gap, not the FLOP count, is what you are fixing.

Tiling, Blocking and the Micro-kernel

The fix has been stable for twenty-five years. Every serious BLAS, from OpenBLAS and BLIS to the GEMM path inside Arm Compute Library, is five nested loops around a hand-written micro-kernel: three levels of blocking, one per level of the hierarchy that offers reuse.

The register tile

The micro-kernel computes an \(M_R \times N_R\) block of \(C\) and keeps it in vector registers for the whole reduction over \(k\). Each \(k\) step loads \(M_R\) values of \(A\) and \(N_R\) values of \(B\) and does \(M_R N_R\) multiply-accumulates, so

for element size \(s\). Operand traffic grows with the tile’s perimeter while work grows with its area, so bigger and squarer is better and the register file is the only limit.

fp32 register tiles on a 32-register NEON machine. Each register holds four fp32 accumulators, so the tile needs \(M_R N_R / 4\) registers for \(C\) alone, plus a few for operands.

8 x 12 is the classic AArch64 fp32 choice: twenty-four registers for \(C\), eight left for operands, intensity 2.4 against 0.25 for the scalar loop. That is 9.6x less operand traffic from register blocking alone, before touching a cache. 16 x 16 needs sixty-four registers and spills, and a spilled accumulator in the innermost loop is worse than the traffic it avoids.

Sizing the cache tiles

Three loops block one dimension each to one cache level, in BLIS names: \(K_C\) the reduction, \(M_C\) the rows, \(N_C\) the columns.The three-level analysis and the MC/NC/KC naming come from the Goto and van de Geijn line of work that became BLIS. The half-capacity heuristic is a starting point, not an optimum: real libraries sweep the final values on the target part, because effective capacity depends on associativity and on what else is resident. The rule: each blocked operand occupies about half its target cache, leaving the rest for streams passing through.

Deriving tile sizes for the illustrative core, fp32, \(M_R = 8\), \(N_R = 12\).

That is a 24 KB B micro-panel in L1, a 256 KB A block filling half of L2 and reused across all 42 micro-kernel calls of the column loop, and a 1.03 MB B block in half the LLC:

for (int jc = 0; jc < N; jc += NC)              // 5: LLC block of B
  for (int pc = 0; pc < K; pc += KC) {          // 4: reduction panel
    pack_B(Bp, B, pc, jc, KC, NC, ldb);         //    -> NR-panel order
    for (int ic = 0; ic < M; ic += MC) {        // 3: L2 block of A
      pack_A(Ap, A, ic, pc, MC, KC, lda);       //    -> MR-panel order
      for (int jr = 0; jr < NC; jr += NR)       // 2: over B micro-panels
        for (int ir = 0; ir < MC; ir += MR)     // 1: over A micro-panels
          micro_kernel(KC, &Ap[ir * KC], &Bp[jr * KC],
                       &C[(ic + ir) * ldc + (jc + jr)], ldc);
    }
  }

Loops 1 and 2 are the only ones touching \(C\); loops 3 to 5 exist purely to keep the right things in the right caches.

Packing, and why the copy pays for itself

Both pack functions copy data that already exists. The first time I saw this I assumed it was a bug.

The micro-kernel wants one 128-bit load to yield four consecutive \(A\) values. In row-major \(A\) the eight rows of a micro-panel are \(K \times 4\) bytes apart: eight streams, eight TLB entries, and on 4 KB pages a block of a wide matrix can touch 128 pages. Packing rewrites it into panel-major order, so the kernel reads one contiguous 256 KB buffer with unit stride and the prefetcher gets it right.

Reuse bounds the cost. A packed \(A\) block serves all \(N_C/N_R\) calls in loop 2, amortising the copy over \(2 M_C K_C N_C\) FLOP, an overhead of \(4/N_C\) bytes per FLOP or 126 FLOP/byte at \(N_C = 504\). Problem 4 does the concrete case: 100 MB moved against 17.2 GFLOP, 1.6 per cent of runtime to remove a 16x penalty. It stops paying only when there is no reuse to amortise, as in a GEMM with \(K\) in the tens.

Vectorising the inner loop

On Armv8.2-A with the dot product extension, SDOT takes two 16-byte vectors, forms sixteen int8 products and accumulates them in groups of four into four int32 lanes: one instruction, sixteen MACs.

#include <arm_neon.h>
#include <stdint.h>

/* Build with -march=armv8.2-a+dotprod. Returns sum a[i]*b[i] in int32. */
int32_t dot_i8(const int8_t *a, const int8_t *b, int n) {
    int32x4_t acc0 = vdupq_n_s32(0);
    int32x4_t acc1 = vdupq_n_s32(0);
    int i = 0;

    for (; i + 32 <= n; i += 32) {              /* two 16-byte chunks */
        int8x16_t a0 = vld1q_s8(a + i);
        int8x16_t b0 = vld1q_s8(b + i);
        int8x16_t a1 = vld1q_s8(a + i + 16);
        int8x16_t b1 = vld1q_s8(b + i + 16);
        acc0 = vdotq_s32(acc0, a0, b0);         /* 16 MACs -> 4 int32 lanes */
        acc1 = vdotq_s32(acc1, a1, b1);
    }

    int32_t sum = vaddvq_s32(vaddq_s32(acc0, acc1));   /* one reduction */
    for (; i < n; ++i) sum += (int32_t)a[i] * (int32_t)b[i];
    return sum;
}

Four things there matter more than the syntax.

Two accumulators, not one. SDOT has a latency of three to four cycles and a throughput of two per cycle, so a single chain serialises at an eighth of peak. Saturating the pipes needs about eight independent chains, and a real micro-kernel has sixteen or more, which is exactly why the register tile is the size it is.

The reduction is outside the loop. vaddvq_s32 is a long-latency cross-lane reduction, and one inside becomes the critical path.

int32 accumulation is safe. The largest int8 product is \(128 \times 127 = 16256\), so overflow needs more than \(2^{31}/16256 \approx 132{,}000\) terms.Comfortable for convolution, where the reduction depth $$C_{in} R S$$ rarely exceeds a few thousand. Not automatically comfortable for a very long GEMM reduction. If you are near the bound, accumulate in int64 or split the reduction, and document it.

A bare dot product is the wrong shape. Four 16-byte loads buy two SDOTs: sixty-four bytes for sixty-four operations, one op per byte, well left of the int8 ridge point of ten, so it is bandwidth bound however well written. A real micro-kernel uses vdotq_laneq_s32, selecting a lane to reuse one loaded \(A\) vector across four \(B\) vectors. A kernel computing one output per reduction can never be fast.

Takeaway: Fast dense kernels are three levels of blocking around a register tile, and the sizes follow from paper arithmetic: intensity \(2 M_R N_R / (s(M_R + N_R))\) subject to the register file, then each cache block at about half its cache. Packing looks wasteful and costs one to two per cent to remove a sixteen-fold penalty.

Four Ways to Compute a Convolution

Unlike GEMM, convolution has several genuinely different algorithms, and picking the wrong one costs more than writing the right one badly.

Notation: input \(C_{in} \times H \times W\), filter \(C_{out} \times C_{in} \times R \times S\), output \(C_{out} \times H_{out} \times W_{out}\), so \(F = 2 \, C_{out} H_{out} W_{out} C_{in} R S\). Running example: \(C_{in} = C_{out} = 64\), \(H = W = 56\), 3x3, stride 1, fp32, which is 231.2 MFLOP, 1.75 MB of compulsory traffic, intensity 132, compute bound at 7.23 ms.

im2col plus GEMM

Materialise every receptive field as a column, giving a \((C_{in} R S) \times (H_{out} W_{out})\) matrix, so the convolution becomes a GEMM with \(M = 64\), \(K = 576\), \(N = 3136\), and \(2 \times 64 \times 576 \times 3136 = 231.2\) MFLOP as before.

The appeal is reusing the GEMM you already tuned; the cost is the buffer. Each input element appears in up to \(RS\) columns, so a 3x3 lowering is nine times the input, \(576 \times 3136 \times 4 = 7.23\) MB against an 803 KB tensor. That is 3.6 times the LLC, so it goes to DRAM and comes back: 14.5 MB, 1.13 ms, on top of a 7.23 ms convolution. Sixteen per cent overhead to reformat data.

Every production library avoids this by lowering one strip of columns at a time, sized for L2, straight into the GEMM’s \(N_C\) loop: a 256-column strip is 590 KB and the overhead vanishes.

Direct convolution

Skip the buffer: a blocked loop nest over \((C_{out}, H_{out}, W_{out}, C_{in}, R, S)\) with a register tile over output channels and pixels, the GEMM micro-kernel with the spatial window folded into the reduction.

Direct wins in four situations: when \(C_{in}\) is small, as in a first layer where the GEMM reduction depth is 147 or less and packing is not amortised; when the convolution is depthwise or grouped and the per-group GEMM degenerates to \(K = RS = 9\); when memory is tight, as on a microcontroller; and when you want to fuse, because the output tile is already in registers at the end of the reduction.

The cost is a specialised kernel per shape and stride, so an unusual shape falls onto a slow generic path. In my experience that is the most common reason one layer runs at a tenth the speed of its neighbours.

Winograd

For small filters you can trade multiplications for additions. Winograd \(F(m \times m, r \times r)\) computes an \(m \times m\) output tile using \((m + r - 1)^2\) element-wise multiplications instead of \(m^2 r^2\). For \(F(2 \times 2, 3 \times 3)\):

The computation is \(Y = A^T [ (G g G^T) \odot (B^T d B) ] A\), with \(g\) the 3x3 filter, \(d\) the 4x4 input tile, \(\odot\) element-wise. For \(F(2,3)\) the matrices \(A\) and \(B\) hold only \(0\) and \(\pm 1\), so input and output transforms are pure additions: about 32 adds per input tile and 24 per output tile.Exact counts depend on the factorisation and on common subexpression elimination; 32 and 24 are the figures usually quoted for $$F(2,3)$$. The filter transform is computed once at model load and costs nothing at inference. In Problem 2 those add 2.7 per cent to the operation count, so on paper the 2.25x survives almost intact.

Not in practice. The element-wise stage becomes 16 small independent GEMMs rather than one large one, the transform passes are bandwidth-bound streams, and the tile buffers are four times the working set of the equivalent direct block. Realistically \(F(2,3)\) delivers 1.3x to 1.8x end to end on a CPU.

Larger tiles are a trap. \(F(4 \times 4, 3 \times 3)\) gives \(144/36 = 4\)x and \(F(6 \times 6, 3 \times 3)\) gives \(324/64 = 5.06\)x, but their transform matrices acquire entries like \(\pm 2, \pm 4, \tfrac{1}{2}, \tfrac{1}{4}\) with fast-growing condition numbers, so the input transform amplifies rounding error through the element-wise stage before the output transform contracts it. In fp32, \(F(4,3)\) is usually acceptable and \(F(6,3)\) marginal; in fp16, \(F(4,3)\) already costs visible accuracy.

Winograd also has hard preconditions: 3x3, stride 1, dilation 1. And it does not compose with int8, because the transforms produce fractional intermediates in a domain where the quantisation scales no longer line up.Published int8 and residue-number Winograd variants do work for specific tile sizes with carefully chosen transforms, but none is a drop-in replacement, which is why production int8 backends fall back to direct or im2col for 3x3.

FFT convolution

Convolution is multiplication in the frequency domain, so you can transform, multiply element-wise, and transform back at \(O(HW \log HW)\) per channel pair instead of \(O(H_{out} W_{out} R S)\). The crossover is where \(\log HW\) beats \(RS\), roughly \(R = S \geq 9\). Modern CNNs use 3x3 and 1x1 almost exclusively, and complex arithmetic doubles storage and makes each multiply four real ones, so FFT wins only in signal-processing front ends and audio models with long 1D kernels.

Choosing between them

Which convolution algorithm to use on a CPU-class edge target.

Takeaway: There is no universally fastest convolution. im2col plus GEMM is the safe default if you tile the lowering buffer; direct wins on small or degenerate channel counts and is the only one that fuses cleanly; Winograd buys 1.3x to 1.8x on 3x3 stride-1 fp32 layers and is unavailable in int8; FFT is for kernels larger than anything in a modern CNN.

Depthwise Separable: Fewer FLOPs, Not Proportionally Less Time

This is the section I would keep if I had to throw the rest away.

A depthwise separable block replaces a standard \(R \times R\) convolution with a depthwise convolution (one filter per channel, no cross-channel mixing) plus a pointwise 1x1 (all the mixing, no spatial extent). For \(C_{in} = C_{out} = C\) the FLOP ratio is \(C R^2 / (R^2 + C)\), tending to 9 for \(R = 3\). That is the arithmetic behind MobileNet, and behind a great many disappointed engineers, because the FLOP reduction and the latency reduction differ by about a factor of two. Take \(C = 64\), \(H = W = 56\), stride 1, fp32.

FLOPs and roofline bounds for a standard block and its separable replacement.

The FLOP reduction is 7.89x, as the formula predicts, and the roofline bound improves by 7.8x. So far the roofline sees no problem.

The problem is that the three kernels do not achieve the same fraction of their bound. The depthwise layer sits at intensity 2.25 against a ridge of 2.5: memory bound, barely. Its 3.61 MFLOP is spread over 64 channels, so each channel is a 56x56 image convolved with nine weights, with no reduction dimension to amortise over and nothing for the register tile to reuse. Every structural advantage the GEMM micro-kernel had is gone. The pointwise layer has a milder opposite problem: as a GEMM it is \(M = 64\), \(K = 64\), \(N = 3136\), so the reduction depth is one eighth of the \(K_C = 512\) we sized for, and \(M < M_C\) so row blocking does nothing.

The same block with illustrative achieved efficiencies. These percentages are not measurements from any product. They are what I would expect from profiling kernels of this shape, and the shape of the result holds even if your numbers differ by twenty per cent.

The speedup is \(11.1 / 2.57 = 4.3\)x against a FLOP reduction of 7.9x. You paid for eight and collected four. The last column says it best: the standard convolution ran at 20.8 GFLOP/s and the block replacing it runs at 11.4, so the separable block does less work and does it at roughly half the efficiency, and the two effects multiply.

This is not an argument against depthwise separable convolutions. Four times faster is four times faster, and I deployed exactly these architectures on robot hardware at Dyson because four times was the difference between running and not running. It is an argument against reading a FLOP count as a latency prediction: it is the first thing I look at and the last thing I trust.

Three things recover part of the gap. Fuse the depthwise into the pointwise, since unfused it writes 803 KB and the pointwise reads it back, 0.125 ms or 45 per cent of the depthwise layer’s runtime. Move to NHWC, so the kernel vectorises across channels with the window walking spatially rather than across space with awkward edges. Quantise to int8, which on a bandwidth-bound layer cuts time by close to four; Problem 1 works that through, and Part 5 covers why depthwise layers need per-channel scales.

Takeaway: Depthwise separable blocks cut FLOPs by up to 9x and latency by typically 3x to 5x, because the depthwise layer is bandwidth bound at intensity near 2 and the pointwise layer is a GEMM too skinny to reach peak. Compare the achieved GFLOP/s of the block you are replacing with the block replacing it; if the second is lower, part of the saving is already spent.

Layout and Fusion

NCHW, NHWC and blocked layouts

A four-dimensional tensor is flattened into a one-dimensional address space, and the order decides which dimension a vector load can reach.

The three layouts you will meet.

NCHW is the historical default because early GPU libraries wanted it. NHWC is what most edge inference wants, because the dominant reduction in a modern network is over channels and NHWC makes it unit-stride. TFLite and XNNPACK are native NHWC, and PyTorch exposes it as channels_last.`to(memory_format=torch.channels_last)` changes strides, not logical shape, so model code is unchanged. On CPU backends with NHWC kernels it is often a 1.5x to 2x win for a one-line change, the highest-value single experiment on the list. Blocked layouts get both properties: NC4HW4 splits the channel axis into groups of 4 and puts the group innermost, giving channels-contiguous vector loads and spatially contiguous planes within a channel block. oneDNN uses nChw8c and nChw16c sized to the vector width, MNN and NCNN use NC4HW4, and Arm Compute Library has its own interleaved GEMM formats.

Hence the heading: the layout the framework hands you is rarely the layout the kernel wants. NCHW boundaries around NHWC kernels mean permutes, and a permute is pure bandwidth with zero arithmetic: reformatting a 56x56x64 fp32 tensor moves 1.6 MB, 0.125 ms, comparable to the entire depthwise convolution it precedes. Choose one layout for the longest run of the graph and reorder only at the boundaries, which is a compiler’s job and much of what Part 4 is about.

Fusing conv, bias and activation

Apply bias then ReLU to our 803 KB convolution output as three kernels: the conv writes 803 KB, bias reads and writes 803 KB, ReLU reads and writes 803 KB, so 4.0 MB and 0.31 ms for two operations doing 2 FLOP per output element.

Fused, both happen in the micro-kernel epilogue while the tile is in registers: one vadd and one vmax per accumulator register against the hundreds of FMAs that produced it, and 803 KB written once for 0.063 ms. That removes 0.25 ms from a 7.2 ms layer, undramatic because the convolution is compute bound. Apply it to the depthwise layer above, where 0.125 ms of 0.28 ms is memory traffic, and it is the whole optimisation.

Batch normalisation is the best case: at inference its statistics are frozen, so it folds into the weights as \(W' = W \gamma / \sqrt{\sigma^2 + \epsilon}\) and \(b' = \beta - \gamma \mu / \sqrt{\sigma^2 + \epsilon}\) and disappears at zero runtime cost. If your converter has not done this, check you exported in eval mode.

What stops fusion

Multiple consumers. If the output also feeds a residual add several layers later, the tensor must exist in memory regardless. You can fuse the activation but you pay the write: residual connections are free in FLOPs and not in bandwidth.

Layout changes. If the producer’s output layout differs from what the consumer needs, the epilogue must scatter rather than store contiguously, and a scattered store in the innermost loop is worse than a separate pass.

Precision changes. An int8 convolution accumulating in int32 and requantising in the epilogue is the standard fused form. But if the next operator has no int8 kernel, the epilogue must dequantise to fp32, four times the bytes, so one unsupported operator can force two conversions and cost more than it saves.

Global reductions. Softmax over a large axis or layer norm needs a statistic over the whole tensor before its first output, so in naive form it cannot fuse into its producer’s epilogue. The online formulations break that dependency and are what makes fused attention possible; Part 8 covers it.

Halos. Pooling fuses into a tile-producing kernel only if its window is handled at tile boundaries, by recomputing a halo or ordering tiles carefully.

Takeaway: Layout and fusion are both bandwidth economics. Pick one layout per subgraph and reorder only at the edges; fuse bias, activation and requantisation into the producer’s register epilogue. Fusion stops at multiple consumers, layout changes, precision changes and global reductions, and each is something you can sometimes change in the model.

Multithreading Without Destroying Locality

Choosing the dimension to split

Splitting each dimension of a GEMM or convolution across threads.

Splitting the reduction is the trap: each thread computes a partial \(M \times N\) result, so you need atomics on every output or \(T\) private buffers of size \(MN\) plus a reduction pass. For a 64x3136 output on four threads that is 3.2 MB of scratch and 3.2 MB of extra reads, to parallelise a reduction of depth 64. It is only right when \(M\) and \(N\) are small and \(K\) huge, as in attention at long context.

Two failure modes destroy locality even with the right dimension. Interleaved rather than blocked partitioning: if thread \(t\) takes channels \(t, t+T, t+2T, \ldots\), no thread has a contiguous slice of the weights and threads sharing an LLC evict each other. False sharing: split the output pixels finer than a cache line and two threads write the same 64-byte line, so every store ping-pongs it between cores. Give each thread a contiguous block, partitioned on an \(N_R\) micro-tile boundary.

Capacity interacts too: \(N_C\) assumed one core owned the 2 MB LLC, but with four threads the share is 512 KB and the tiles should be re-derived. The symptom of missing this is a kernel that scales beautifully to two threads and then stops.

The cost of the barrier

Every layer boundary is a synchronisation point. A spin barrier across four cores on one cluster is a handful of cache line transfers, call it 2 us with imbalance. Waking four threads sleeping on a futex costs a system call, a scheduler decision, possibly a migration and a clock ramp: 20 to 50 us is normal.Order-of-magnitude figures, not measurements from a specific part. Get your own by running a network of empty kernels through the real thread pool and reading off the per-layer floor: that floor is your barrier cost, one of the most useful numbers to have when reasoning about batch-1 latency. At batch 1 that dominates, because per-layer work is small.

Per-layer latency and parallel efficiency for a 54-layer network at batch 1, with 180 us of serial work per layer.

Same kernels, same hardware, same partitioning: the only difference is pool policy, worth 1.24 ms of 3.78. Shrink the layers to 40 us serial and four cores give no speedup at all under the sleeping pool, \(40/4 + 25 = 35\) us against 40 single-threaded. Use a persistent pool that spins for a few microseconds before sleeping; finding that knob is often the highest-return change available to a batch-1 workload. The other lever is fewer barriers, which is another argument for fusion.

Takeaway: Split the output dimensions in contiguous blocks, never the reduction, and never finer than a cache line. At batch 1 the per-layer synchronisation cost is comparable to the layer itself, so a spinning thread pool and fewer, fatter fused kernels beat any amount of inner-loop tuning.

Knowing When You Are Done

Optimisation without a stopping criterion expands to fill the time available. Mine is two numbers: measure the time, then compute both

with \(F\) the analytically counted operations and \(B\) the compulsory traffic. The larger fraction names your wall. If neither exceeds about forty per cent, you have a third problem: instruction overhead, a latency chain, a bad layout, thread imbalance, or a shape off the fast path.

Get the peaks by measurement: for compute, a loop of independent FMAs on register-resident data with enough accumulator chains to cover the latency; for bandwidth, a STREAM triad sized several times the last-level cache. Datasheet numbers assume all cores at maximum clock with no thermal limit, which is not the condition your kernel runs in.

What a good efficiency number looks like on an edge CPU core. Rules of thumb from experience, not guarantees.

Count FLOPs analytically from the shapes, not from a hardware counter, which includes address arithmetic and predicated lanes and will flatter you. And with a reduced-arithmetic algorithm like Winograd, report the direct FLOP count over your time.Otherwise a Winograd kernel that is genuinely 1.5x faster reports an efficiency 2.25x lower than it deserves. Quoting effective GFLOP/s against the direct operation count occasionally exceeds the machine's peak, which is fine: it means the algorithm won.

Finally, the Amdahl check, which overrides all of the above. Taking a layer that is three per cent of runtime from 40 to 80 per cent efficiency buys 1.5 per cent overall; taking the layer that is forty per cent from 40 to 60 buys thirteen. The criterion is not “this kernel is as fast as it can be” but “the next hour is better spent somewhere else”.

A Few Problems to Work

Problem 1. A depthwise 3x3 stride-1 convolution has \(H = W = 112\) and \(C = 32\), on a core with 40 GFLOP/s fp32 peak, 160 GOP/s int8 peak and 15 GB/s of bandwidth. Compute the operation count, compulsory traffic, intensity and roofline bound in fp32; then repeat in int8 and say where the speedup comes from.

Problem 2. A 3x3 stride-1 convolution has \(C_{in} = C_{out} = 128\) and \(H = W = 28\) with padding preserving size. Count the multiplications for direct convolution and for Winograd \(F(2 \times 2, 3 \times 3)\), add transforms at 32 additions per input tile per input channel and 24 per output tile per output channel, give the net operation-count speedup, and say why the measured speedup is lower.

Problem 3. A 54-layer network averages 180 us of serial work per layer on 4 cores. Compare total latency and parallel efficiency for a spinning pool with a 2 us barrier and a sleeping pool with a 25 us barrier. Then: how many layer pairs would you fuse to recover half the difference?

Problem 4. For a 2048 x 2048 x 2048 fp32 GEMM with \(M_C = 256\), \(K_C = 512\), \(N_C = 1024\), count total packing traffic and express it as a percentage of runtime, assuming 70 per cent of a 50 GFLOP/s peak and packing traffic at 12.8 GB/s.

What’s Next

We now have the inside of an operator: why the obvious loop reaches six per cent of peak, what the three levels of blocking are for, how to choose a convolution algorithm, and why a FLOP saving is not a latency saving. What we did by hand is what a compiler is supposed to do for us: choose layouts, decide fusion boundaries, pick tile sizes, select algorithms per shape, schedule the graph. Part 4 is about that machinery, and about the gap between what a compiler can do in principle and what it does to your model on a Tuesday afternoon.

That’s all for Part 3! For Part 4, on compilers, graphs and runtimes, click here.