Inside an Edge Accelerator

Part 1 gave you a model of a machine: a peak compute rate, a peak memory bandwidth, and a ridge point where the two cross. That model is only worth anything if you can fill in the numbers for a device you actually have to ship on. This chapter is about where those numbers come from, which means it is about silicon: what instructions the CPU really has, how the memory hierarchy is shaped, what an NPU is underneath its headline TOPS figure, and why the chip that gave you 28 ms per frame on Monday morning gives you 44 ms on Tuesday afternoon.

I work on this boundary for a living. At Arm I write and optimise ML libraries and compiler back ends, which mostly means staring at instruction schedules and cache behaviour until a kernel stops being stupid. Before that I spent two years at Dyson putting convolutional networks onto robot hardware, where the model that wins the accuracy bake-off and the model that survives a hot plastic enclosure are rarely the same one. Both jobs taught me that you can predict most of a deployment’s performance from a datasheet and an afternoon of arithmetic, and that almost nobody does it.

Four Ways to Run a Model Outside a Datacentre

Open the block diagram of any phone, camera, car module or vacuum robot. You will find at most four kinds of compute that can usefully run a neural network, and usually three of them on a single die.

The four classes of edge compute, with order-of-magnitude envelopes. These are class-level ranges, not any specific product; read them as “what to expect before you look up the real part”.

The CPU and its vector units

The CPU is the only engine guaranteed to run your model: every operator, every shape, all your control flow, and a profiler you already know. Its vector units are genuinely fast, 64 multiply-accumulates per cycle on a big core with the int8 matrix extension. What it is bad at is energy: it pays for an out-of-order pipeline, branch prediction, address translation and coherent caches on every instruction, and a matrix multiply needs none of that. On a large dense convolution it is 5 to 20 times worse in operations per joule than a fixed-function MAC array on the same node.

The integrated GPU

The integrated GPU is a wide throughput machine that already exists on the die because something has to draw the screen. It is excellent at fp16 and at anything image-shaped: resize, colour conversion, warping, non-maximum suppression. Its weakness is fixed cost per dispatch: a kernel launch, a queue submission and a fence wait can cost 50 to 200 microseconds, so a graph of forty small operators spends longer in the driver than in arithmetic. It also shares bandwidth and power with whatever composites the user interface.

NPUs and DSPs

An NPU is a fixed-function matrix engine with its own local SRAM and command stream. It is the best operations-per-joule option on the die, often by an order of magnitude over the CPU, and that is the entire reason it exists. The price is rigidity, which gets a section of its own below.

A vector DSP sits between the two: a programmable core with a very wide SIMD unit, often 512 or 1024 bits, tuned for integer signal processing rather than floating point.The best known family in phones pairs 1024-bit vectors with a software-managed local memory rather than a data cache. That matters more than the width: the compiler, not the hardware, decides what is resident. DSPs own the always-on tier, where a keyword spotter or a presence detector must run at a few milliwatts while the rest of the SoC sleeps.

Small discrete accelerators

If you cannot change the host, you bolt an accelerator onto it: the M.2, USB and small PCIe category. A Google Coral Edge TPU is publicly specified at 4 Top/s int8 for roughly 2 W, and Jetson-class modules reach tens of Top/s in a 7 to 60 W envelope. Both are vendor-published peaks, reachable on a synthetic kernel with everything resident.

The trap is the link. A USB 3 accelerator that infers in 3 ms is not a 3 ms solution if moving a 640x640x3 tensor across the bus and waiting for the completion interrupt costs another 4 ms. Measure the round trip from frame in to result out, never the device-side kernel time.

Takeaway: the CPU always works, the GPU is wide but has expensive dispatches, the NPU has the best operations per joule and the least flexibility, and the DSP owns the always-on tier. Most products use two or three in one pipeline.

Arm CPUs in Detail

This is the part I know best, so we will go slowly and do the arithmetic properly.

NEON: 128 bits, fixed

NEON (formally Advanced SIMD) gives you 32 architectural registers of exactly 128 bits, addressed by lane layout: v0.16b is sixteen 8-bit lanes, v0.8h eight 16-bit lanes, v0.4s four 32-bit lanes. The width is fixed by the architecture, so a NEON binary compiled in 2013 runs unchanged on a 2026 core.

Fixed width costs you tails. With a reduction length of 1000 and 16 bytes per vector you do 62 full iterations and then a scalar loop for the last 8 elements: often a third of the source of a hand-written kernel, contributing nothing to throughput.

SVE and SVE2: vectors of unknown length

The Scalable Vector Extension takes the opposite approach. Vector length is an implementation choice, any multiple of 128 bits from 128 up to 2048, and the program does not know it at compile time. You never write “16 bytes”; you write svcntb(), which returns the byte count of a vector on whatever core you are running on.

Two features make this practical. Predication: 16 predicate registers let every load, store and arithmetic instruction be masked per lane, and the whilelt idiom builds a predicate true for lanes below the loop bound, so the last partial iteration runs the same instructions as every other one. Length-agnostic idioms: loop increments use INCB/CNTB rather than an immediate, and reductions like svaddv fold a whole vector whatever its width.

Here is a complete int8 dot product with no tail loop:

#include <arm_sve.h>
#include <stdint.h>

int32_t sve_dot_i8(const int8_t *a, const int8_t *b, int64_t n) {
    svint32_t acc = svdup_s32(0);
    int64_t i = 0;
    svbool_t pg = svwhilelt_b8(i, n);

    while (svptest_any(svptrue_b8(), pg)) {
        svint8_t va = svld1_s8(pg, a + i);   /* inactive lanes read as 0 */
        svint8_t vb = svld1_s8(pg, b + i);
        acc = svdot_s32(acc, va, vb);        /* 4-way dot into int32 lanes */
        i += svcntb();                       /* vector length in bytes */
        pg = svwhilelt_b8(i, n);
    }
    return svaddv_s32(svptrue_b32(), acc);
}

That is correct for n = 3 and for n = 3000000, on a 128-bit implementation and a 512-bit one, from a single binary.`svdot_s32` is unpredicated, which looks like a tail bug until you notice that `svld1_s8` under a predicate returns zero in inactive lanes, and zero times anything is zero. Without predicated loads you would need an explicit mask-and-clear, which is what makes hand-written tail code buggy.

SVE2 adds the integer and fixed-point operations DSP and media code needs, so it is a general replacement for NEON rather than an HPC extension. Length agnosticism matters because you ship one binary to a fleet containing several cores. Without it you compile N versions and dispatch at runtime, or compile for the narrowest implementation and waste half the throughput on the wide ones.

SDOT, SMMLA and the int8 step change

Before Armv8.2, int8 inference on an Arm CPU was awkward. With no int8 multiply-accumulate, the pattern was widen-multiply-accumulate: SMULL/SMLAL give 16-bit products from 8-bit inputs, and you widened again to 32 bits before the 16-bit accumulator overflowed. Much bookkeeping, and int8 was barely faster than fp32. The dot product instructions changed that.`SDOT`/`UDOT` are FEAT_DotProd, optional from Armv8.2-A and mandatory from Armv8.4-A; `SMMLA`/`UMMLA`/`USMMLA` are FEAT_I8MM and `BFDOT`/`BFMMLA` are FEAT_BF16, both mandatory from Armv8.6-A. Detect them at runtime from the auxiliary vector (`HWCAP_ASIMDDP`, `HWCAP2_I8MM`, `HWCAP2_BF16`), because the compile-time target rarely matches the fleet.

SDOT Vd.4S, Vn.16B, Vm.16B treats each source as four groups of four signed bytes. For each of the four 32-bit destination lanes it multiplies the corresponding group element-wise and adds all four products into that lane’s existing value. Sixteen multiplies and sixteen additions in one instruction, accumulating straight into 32 bits so overflow is a non-issue. Take:

Vn.16B = [ 1,  2,  3,  4 | 5,  6,  7,  8 | -1, -2, -3, -4 |  2, 2,  2, 2]
Vm.16B = [ 1,  1,  1,  1 | 2,  0,  2,  0 |  3,  3,  3,  3 | -1, 1, -1, 1]
Vd.4S  = [ 0, 0, 0, 0 ]

After SDOT Vd.4S, Vn.16B, Vm.16B:

so Vd.4S = [10, 24, -30, 0]: a length-4 reduction in each of four lanes.

SMMLA does a small matrix product instead. It reads Vn as a 2x8 int8 matrix, Vm as the transpose of a 2x8 (so, an 8x2), multiplies them and accumulates the 2x2 int32 result into Vd.4S. That is \(2 \times 2 \times 8 = 32\) multiply-accumulates per instruction, double SDOT, from the same 256 bits of input, because each input byte now feeds two products instead of one. In C you reach both through ACLE intrinsics:

#include <arm_neon.h>

int32x4_t acc = vdupq_n_s32(0);
int8x16_t a = vld1q_s8(pa), b = vld1q_s8(pb);

acc = vdotq_s32(acc, a, b);    /* SDOT:  16 MACs, 4 independent lanes */
acc = vmmlaq_s32(acc, a, b);   /* SMMLA: 32 MACs, a 2x2 int32 tile    */

Multiply-accumulates performed by one 128-bit instruction, by data type. This is the most useful table on the page when estimating a CPU’s peak.

Now the peak. Let a big core run at 2.0 GHz with two 128-bit SIMD pipes that can each issue SMMLA every cycle.Issue width varies: small in-order cores typically have one SIMD pipe, big out-of-order cores two to four, and not all pipes execute every instruction. Check the core's software optimisation guide for the specific instruction's throughput; assuming otherwise is the commonest source of wrong peak estimates I see.

Four such cores give 1.02 Top/s of int8 before you touch the NPU. With SDOT you get half that; with the old SMLAL path about a quarter. This is why int8 quantisation (Part 5) and feature detection must be designed together: quantising buys you nothing if your runtime dispatches to a kernel written for a core from 2016.

Takeaway: one SDOT is 16 int8 MACs and one SMMLA is 32. Peak int8 per core is SIMD pipes times MACs per instruction times two times clock, which lands near 256 Gop/s on a modern big core at 2 GHz.

The Memory Hierarchy and Why Kernels Are Tiled

Having 64 MACs per cycle is meaningless if you cannot deliver operands at 64 MACs per cycle. Everything interesting about kernel engineering follows from that sentence.

Representative numbers

Approximate latency, bandwidth and relative energy for each level of a typical Arm SoC memory hierarchy. These are orders of magnitude for a mid-range mobile-class part, not the specification of any product. Look yours up; use these to check that you looked it up correctly.

The latency column spans more than two orders of magnitude, and out-of-order execution hides only the top of it. The energy column is why the NPU exists: at 45 nm, Horowitz’s widely cited figures put an 8-bit integer multiply at about 0.2 pJ and a 32-bit DRAM read at about 640 pJ, a ratio near 3000 to 1.M. Horowitz, "Computing's Energy Problem (and what we can do about it)", ISSCC 2014. The absolute numbers shrink on newer nodes but the ratios hold, because arithmetic scales with the process and off-chip I/O largely does not. Arithmetic is nearly free. Moving data is what you pay for.

Where tiling comes from

Consider \(C = AB\) with \(A\) of shape \(M \times K\), \(B\) of shape \(K \times N\), int8 inputs. The naive triple loop reads \(K\) bytes from \(A\) and \(K\) from \(B\) for every one of the \(MN\) outputs, so it moves \(2MNK\) bytes and does \(MNK\) MACs. That is 0.5 MACs per byte, forever, whatever the matrix size.

Our core wants 64 MACs per cycle, so each level imposes a required reuse: the MACs per byte you must achieve against it to avoid being limited by it. Take a 2.0 GHz core and four big cores sharing an LPDDR5-6400 64-bit interface (51.2 GB/s peak, so 12.8 GB/s per core, which is 6.4 B/cycle):

Two levels of tiling supply those numbers. Register tiling holds an \(M_r \times N_r\) block of \(C\) in vector registers; each step along \(K\) loads \(M_r + N_r\) bytes and does \(M_r N_r\) MACs, so reuse against L1 is \(M_r N_r / (M_r + N_r)\). An 8x12 tile gives \(96/20 = 4.8\), above the required 2, and its 96 int32 accumulators occupy exactly 24 of the 32 NEON registers. Real micro-kernels are 8x12 or 12x8 for that reason: the tile is chosen by the register file, not by taste.

Cache blocking picks an \(M_c \times N_c\) block of \(C\) and sweeps the full \(K\) dimension through it, giving DRAM reuse \(M_c N_c / (M_c + N_c)\). With \(M_c = N_c = 256\) that is \(65536/512 = 128\) MACs per byte, twelve times the required 10; even a 64x64 block gives 32.

A tiled kernel exists, then, because the required reuse at each level is fixed by the ratio of compute throughput to that level’s bandwidth, and only a blocked loop nest supplies it. Real implementations also pack each panel contiguously first, trading a little bandwidth to turn strided loads into streams the prefetcher understands. We build one in Part 3.

Takeaway: every cache level sets a required reuse equal to MACs per cycle divided by bytes per cycle. Register tiling supplies the L1 number and cache blocking supplies the DRAM number; a kernel missing either runs at a fraction of peak however good its inner loop is.

Inside an NPU

The MAC array and its dataflow

Strip away the marketing and an NPU is a two-dimensional grid of multiply-accumulate units, a local SRAM, a DMA engine and a sequencer walking a pre-compiled command stream. The grid is the point: an \(R \times C\) array does \(RC\) MACs per cycle from one instruction, because control logic is amortised across thousands of multipliers instead of a handful of SIMD lanes. A 128x128 array at 1 GHz does 16,384 MACs per cycle, 16.4 TMAC/s, which a datasheet prints as 32.8 TOPS.

How operands move through the grid is the dataflow. Weight stationary holds one weight in each processing element while activations stream in and partial sums propagate across. This is the classic systolic array: excellent when each weight is reused many times, bad when it is not, since a fully connected layer at batch 1 uses each weight once and the array spends its time being reloaded. Output stationary gives each element one accumulator in a local register while weights and activations both stream, avoiding partial-sum movement (int32, wider than the int8 inputs) and suiting deep reductions.

Either way, utilisation decides your real throughput. If the array maps input channels to rows and output channels to columns, a layer narrower than the array leaves rows or columns idle.

Worked example. A 128x128 output-stationary array at 1 GHz, peak 16.4 TMAC/s, running a layer with 48 input channels and 320 output channels.

• Rows: 48 of 128, so 37.5%.
• Columns: 320 needs \(\lceil 320/128 \rceil = 3\) passes, using \(320/(3 \times 128) = 83.3\%\).
• Combined: \(0.375 \times 0.833 = 31.25\%\), giving \(16.4 \times 0.3125 = 5.1\) TMAC/s.

The “32.8 TOPS” chip just delivered 10.2 TOPS on a perfectly ordinary layer, and nothing was broken. It is also why depthwise separable convolutions, which every efficient mobile architecture is full of, disappoint on systolic hardware: a depthwise convolution has no input-channel reduction, so it uses one row.Good NPU compilers map them differently, for example putting spatial positions on the reduction axis, or fusing the depthwise and following pointwise convolution. The ceiling is still far below a dense convolution, which is why MAC count and latency correlate so poorly on NPUs.

On-chip SRAM is the real constraint

The array is rarely the bottleneck; the local SRAM is. Typical on-NPU SRAM is 1 to 8 MB, and per tile of work it holds the input tile, the layer’s weights, the output tile and usually a double buffer so DMA overlaps compute. If the working set does not fit the compiler tiles it, and tiling re-reads from DRAM. If the weights alone do not fit, you re-read them every inference.

Worked example. A 2048x2048 int8 fully connected layer has 4,194,304 weights (4 MiB) and does 4,194,304 MACs at batch 1. If they stream from a 51.2 GB/s interface:

The layer is 320 times memory bound and the array runs at 0.3% utilisation. Nothing can be done at batch 1, because the arithmetic intensity of a matrix-vector product is fixed near 1 MAC per byte. This is the most important fact about running transformers on edge NPUs, and we return to it in Part 8.

Dynamic shapes, odd operators and the fallback problem

NPU compilers do ahead-of-time memory planning: they fix tile sizes, allocate SRAM addresses, schedule DMA and emit a static command stream before the first inference. That is how they get their efficiency, and the source of every frustration you will have with them.

Dynamic shapes break the plan. If sequence length or detection count is unknown at compile time, the compiler must recompile per shape (hundreds of milliseconds, unusable per frame), pad to a maximum shape (wasting proportional compute), or bucket into a few fixed shapes. There is no fourth option, and bucketing is what I would do.

Unsupported operators break the graph. The operator set is hardware plus a limited programmable layer, so a custom activation, an unusual reduction axis or a five-dimensional transpose can all miss. The runtime then runs that piece on the CPU, and the cost is not the cost of the operator. One miss in the middle turns a network into three subgraphs: NPU, CPU, NPU. Each boundary costs a driver synchronisation, a layout conversion (NPU-internal tiled formats are not NHWC), and cache maintenance over the intermediate tensor in both directions. At a few hundred microseconds a crossing on a 5 ms network, one operator takes 10 to 20% of your frame budget while doing almost no arithmetic.

At Dyson this was comfortably the commonest reason a model that looked fine in the framework disappointed on the robot. The fix is to change the model, not fight the compiler: rebuild the exotic operator from supported ones, move post-processing out of the graph, and freeze your shapes. Three subgraphs down to one beats any amount of kernel tuning.

Takeaway: an NPU’s real throughput is peak times array utilisation, and utilisation collapses on narrow, depthwise and batch-1 layers. One unsupported operator can cost far more than the operator does, because it buys you two extra subgraph boundaries.

Shared Memory Is Not Free Memory

On an SoC the CPU, GPU and NPU usually share the same DRAM, so tensor handoff looks free and everyone says “zero copy”. What is genuinely free is the memcpy you avoid: on a 4 MB tensor at 10 GB/s effective, about 0.8 ms. Four things are not free.

Cache maintenance. If the accelerator is not I/O coherent, the CPU must write back dirty lines before the accelerator reads the buffer and invalidate stale lines before reading its output, per cache line, over the whole buffer.Cache lines are 64 bytes on essentially all current Arm application cores, but read it from `CTR_EL0` rather than assuming: instruction and data line sizes can differ, and some cores use 128. A 4 MB buffer is 65,536 lines, so a clean plus an invalidate is 131,072 operations; at one line every three to four cycles that is roughly 0.2 to 0.25 ms at 2 GHz, plus writeback bandwidth. That figure is illustrative, and the point is that it is the same order as the copy you avoided.

Bandwidth contention. Zero copy does not create bandwidth. When the NPU streams weights at 30 GB/s from a 51.2 GB/s interface, the CPU thread doing your letterbox resize gets what is left and can run at half speed. I have seen post-processing blamed for a regression whose real cause was that the NPU work ahead of it had got faster and was now saturating the memory controller.

Layout mismatch. If the producer writes NHWC fp32 with arbitrary stride and the consumer wants a tiled int8 layout with 64-byte alignment, “zero copy” silently becomes a repack: a copy with a transpose in it. Allocate in the layout the driver wants.

Import cost. Mapping an external buffer into the device address space, and possibly the IOMMU, costs tens to hundreds of microseconds. Do it at model load; a pipeline that re-imports every frame can spend longer in the driver than in the network.

Clocks, Heat and Honest Benchmarking

Why the first run is the fastest

Almost. The honest benchmark curve has three phases.

Phase 1, the first few iterations: slow. Cold caches, first-touch page faults on lazily allocated buffers, the runtime choosing kernels, the allocator growing arenas. The first inference is often 2 to 10 times steady state, hence warm-up iterations.

Phase 2, roughly iteration 5 to a few hundred: fastest. Caches hot, and the governor has raised the clock, often to a boost frequency only available while the die is cool. This is the number that ends up in marketing material.

Phase 3, after the thermal mass saturates: the truth. Dynamic voltage and frequency scaling responds to junction temperature, to a power budget shared across the SoC, and on enclosed devices to a skin temperature limit that often binds before anything on the die is hot. The clock drops and throughput settles 20 to 50% below phase 2. How long that takes depends on thermal mass: tens of seconds for a phone, minutes for a metal chassis, never on a heatsink. A robot vacuum is a plastic box with a motor in it, so at Dyson the sustained number was the only one that mattered.

To measure honestly:

1. Warm up and discard at least 20 iterations, or 2 seconds, whichever is longer.
2. Run for a fixed wall-clock duration, not a fixed iteration count. Sixty seconds minimum. A fixed count reports phase 2 for fast models and phase 3 for slow ones, so your comparisons are not comparisons.
3. Log clock frequency and temperature (scaling_cur_freq and the thermal zones). Without them you cannot tell a throttled run from a slow kernel.
4. Pin clocks where you can: with root, set the governor to performance, or pin scaling_min_freq and scaling_max_freq together. That makes a measurement repeatable, which is what you want while optimising and not what you want for an acceptance test.
5. Report both numbers with your real duty cycle. “28 ms burst, 44 ms sustained after 25 s at 45 C ambient, deployed at 28% duty” is honest; “28 ms” is not. At a low duty cycle the device may never reach phase 3, so the sustained figure is pessimistic, but say so.

big.LITTLE and why you pin the inference thread

Arm SoCs are heterogeneous: large out-of-order cores at a high clock alongside small efficient cores at a lower one, sharing a coherent interconnect. Per-core int8 throughput commonly differs by 3 to 5 times, because the big core has both a higher clock and more SIMD pipes.

Two consequences cost real time. Static work splitting is a disaster: split a layer into equal chunks across all 8 cores and the big cores idle while the little cores grind, so the operation runs at the speed of the slowest core. Thread migration destroys your working set: the scheduler places threads from a utilisation signal that ramps over tens of milliseconds, so a freshly woken inference thread can start on a little core, be promoted mid-inference and arrive with cold L1 and L2. For a 5 ms inference the ramp never finishes.

The remedies, in order of effort:

• Pin. sched_setaffinity to the big cluster, or a cpuset. ONNX Runtime, TFLite and the Arm Compute Library all expose thread affinity. Do this first and measure again.
• Size the pool to the cluster, not the chip. Four threads on four big cores beats eight spread across everything for a latency-sensitive inference.
• Use dynamic work stealing if you must use all cores. A queue where each core takes the next tile when free extracts real throughput from heterogeneity. Problem 3 works the arithmetic.

Concentrating work on fewer cores raises their clock but also concentrates power in a small area of die, so pinning is normally faster and often less efficient than spreading. Which matters depends on whether your constraint is latency or battery.

Takeaway: the fastest run is never the honest one. Warm up, run for a fixed duration, log clocks and temperature, pin the inference thread to the big cluster, and report burst and sustained numbers side by side with your real duty cycle.

Reading a Datasheet for Three Numbers

Everything above exists so you can do this quickly. From Part 1 you need three numbers: peak compute at your precision, peak memory bandwidth, and the ridge point where they cross.

How to translate common datasheet lines into roofline inputs.

Worked example: a hypothetical Device X. The datasheet gives 4 big cores at 2.6 GHz with FEAT_I8MM and two SIMD pipes, 4 little cores at 2.0 GHz with FEAT_DotProd and one SIMD pipe, an 8 TOPS int8 NPU, and 4 x 16-bit LPDDR5-6400.

Peak compute. Big cores: \(2 \times 32 \times 2 = 128\) ops/cycle each, so \(4 \times 128 \times 2.6\times10^{9} = 1.331\) Top/s. Little cores: \(1 \times 16 \times 2 = 32\) ops/cycle each, so \(4 \times 32 \times 2.0\times10^{9} = 0.256\) Top/s. CPU total 1.59 Top/s; the NPU’s 8 TOPS is 4 TMAC/s. Two rooflines, and you should draw both.

Peak bandwidth. \(6400 \times 8\ \text{B} = 51.2\) GB/s, derated to 70% for refresh, bank conflicts and read/write turnaround: 35.8 GB/s.60 to 80% is the usual range for a streaming pattern on LPDDR; random access across many pages does far worse, sometimes 20%. A five-line STREAM-style triad on the target gives the real number in a minute.

Ridge point. CPU: \(1.59\times10^{12} / 35.8\times10^{9} = 44\) ops per byte, or 22 MACs per byte. NPU: \(8\times10^{12} / 35.8\times10^{9} = 223\) ops per byte, or 112 MACs per byte.

Sit with that second number. The NPU has five times the compute of the CPU cluster and the same memory system, so its ridge point is five times further right. A model at 40 MACs per byte is comfortably compute bound on the CPU and badly memory bound on the NPU. Moving it saves energy, because the NPU burns far less per operation, but it need not make it faster. This arithmetic is why people are surprised that “the NPU is not helping”.

Takeaway: peak compute is cores times pipes times MACs per instruction times two times clock; peak bandwidth is data rate times bus width over eight, derated; the ridge point is their ratio. Compute one roofline per engine, because the faster engine has the harder ridge point to clear.

A Few Problems to Work

Problem 1. A device has 6 identical cores at 2.4 GHz, each with two 128-bit SIMD pipes and FEAT_DotProd but not FEAT_I8MM. Memory is a single 32-bit LPDDR4X-4266 channel. (a) Peak int8 throughput? (b) With FEAT_I8MM? (c) The ridge point in the I8MM case? (d) A 3x3 depthwise convolution with 32 channels on a 112x112 int8 feature map, stride 1, same padding: compute or memory bound, and by how much?

Problem 2. An NPU has a 128x128 output-stationary MAC array at 1.0 GHz, 2 MB of local SRAM, and shares a 51.2 GB/s LPDDR5 interface. (a) Peak int8 in TOPS? (b) A convolution has 96 input channels and 200 output channels: array utilisation and effective throughput, with input channels on rows and output channels on columns? (c) A 1024x1024 int8 fully connected layer at batch 1 streaming its weights from DRAM: what array utilisation?

Problem 3. A layer is decomposed into 32 independent tiles. A big core completes a tile in 4 ms, a little core in 13 ms. The device has 4 big and 4 little cores. Find the makespan for (a) a static even split of 8 tiles per core, (b) the 4 big cores only, (c) a dynamic work-stealing queue over all 8 cores.

Problem 4. A detector runs at 28 ms per frame for the first 25 seconds of a continuous run, then throttles to a steady 44 ms. (a) What mean frame rate does a 5 minute continuous benchmark report? (b) A 10 second benchmark? (c) The product needs 15 fps. Assuming the device holds its burst clock indefinitely at duty cycles at or below 50%, is 15 fps sustainable, and what is the maximum sustainable frame rate?

What’s Next

We now have the machine: the instructions that do the arithmetic, the hierarchy that feeds them, the accelerators alongside, and the thermal and scheduling behaviour that decides which measurements were real. What we have not done is write a fast kernel. That is next: taking a matrix multiply from the naive triple loop to something that reaches 70% of the peak we just computed, one blocking decision at a time.

That’s all for Part 2! For Part 3, on kernels and where the time actually goes, click here.