One level down: #004 the same compiler on Apple's chips

The previous article ended on a question: does designing a compiler for an LLM inference engine work on Apple's chips? Because in this particular case, there is no question of writing some kind of assembly, as I did for CPUs, for Vulkan with SPIR-V or for NVIDIA GPUs with PTX.

Here, the only option is Metal, Apple's API, on unified-memory machines (memory shared between the processor and the GPU). Metal has its own language, the Metal Shading Language: a source language, close to C++, that the driver compiles itself. Apple documents neither the instruction set of its GPUs nor the intermediate form its compiler produces.

The information does partly exist, thanks to Asahi Linux, an incredible reverse-engineering project: its developers documented a good part of Apple's GPU to write free drivers and run Linux on the first generations of Apple Silicon Macs. But nothing usable for our needs. So here, Herbert does not write assembly: it writes source text at startup and hands it to the driver.

And that is not all. I worked with three Apple machines. Yet they have almost nothing in common.

In short

• On a Mac Studio M3 Ultra, Herbert produces the same tokens as mlx-lm, up to ties. Up to 8,192 tokens, it is ahead of mlx-lm and llama.cpp at prefill on all six models, by 4 to 31% over mlx-lm (which did not load the 27B). It is ahead at decode everywhere, by 9 to 34% over the faster of the two; the lead holds after a 32,768-token prompt.
• In the first campaigns of this new Herbert engine, llama.cpp was ahead at prefill on all six models. Since then, I have carefully adapted Herbert to measure several shapes of each kernel at load time and keep the fastest. For a while, it picked faster kernels... but wrong ones. After yet another fix, Herbert checks each candidate before taking it.
• Of course, on long prompts, Herbert got wrecked: at 81,920 tokens, 1,128 tokens per second against 4,848 for mlx-lm. The prefill attention had stayed in scalar arithmetic :-/ Rewritten on the chip's matrix units, it gives 4,827, level with mlx-lm (4,784 in the same campaign). The lead melts away with length: at 61,440 tokens, the 2B is 1% behind.

1. A Mac Studio M3 Ultra, the big one of the bunch: 96 GB of memory, 819 GB/s on paper to read it, enough to hold in bf16 models that no consumer card can load. With it, the point is to waste nothing: a GPU this size is quickly left waiting for its data. Well, I have a lot to say about the memory controller here, because the memory bandwidth is not the one advertised, and I was rather surprised to see that the CPU did not have full access to it. I will probably write a post or an article about it.

2. A MacBook Air M1, my oldest device: 8 GB, around 60 GB/s measured, and so unable to hold anything bigger than a 2B. It was Apple's beginning in computer processors, and it shows: its processor's matrix unit does not know the bf16 number format, the format of the whole series, and its GPU computes it slowly: llama.cpp itself runs 18% faster there in f16 than in bf16. I will probably come back to this in more detail in a coming post, maybe even an article.

3. A MacBook Air M5, the newcomer: 32 GB and matrix compute units the other two lack. Apple calls them "Neural Accelerators" and puts one in each GPU core; you reach them through Metal 4's tensor operations, such as matmul2d. On this chip, they multiply matrices two to three times faster than the classic matrix units. On the other hand, MacBook Air oblige: no fan. So it runs very fast for thirty seconds, then it heats up and no longer measures the same thing: that is throttling. Past a certain temperature, the system lowers the chip's clock to protect it: everything slows down.

The idea was to write a single Herbert that adapts itself to all three: at load time, it looks at which chip it runs on, measures what it can do and adjusts to it. That is the heart of my work, but this time in Metal. The kernels, the tiles, the block sizes are chosen on the machine, not in advance; the number format, bf16 or f16, is chosen at build time, depending on what the chip can do. That is what this whole article is about.

The source text Herbert writes is made for the shape of the model and for the chip: every known dimension appears in it as a constant. The driver compiles it without "fast math", the mode where the compiler allows itself to reorder computations and take rounding shortcuts: the basic operations (addition, multiplication, fused multiply-add) are rounded as the IEEE standard prescribes. No .metal file written in advance, no MLX, no Metal Performance Shaders, Apple's compute library: the same emitter as under Vulkan, with the same numerical contract (weights, activations, K/V cache and recurrence states in bf16, accumulators in f32). And Herbert rounds only where Hugging Face's reference code rounds too.

The Mac Studio's memory also allows a wider choice of models. On the 4080 and its 16 GB, I stopped at the 4B. Here, six Qwen models fit in bf16: Qwen3.5 0.8B, 2B, 4B and 9B, Qwen3.8-27B (a 55 GB dense model) and Qwen3.6-35B-A3B, a mixture-of-experts (MoE) model of which only 3 billion parameters are said to be "active" and used for each token. I also ported Gemma 4, the model everything started with (more on that just below), but its figures will come later.

Months of work

This engine is not the first one I have written for Apple's chips. It is actually on them that I started writing GPU kernels, months before Vulkan and CUDA. And it was the hardest work, if I set aside what I am doing on Blackhole, Tenstorrent's card, which I will talk about when I am done.

I first wrote my matrix multiplications by hand, in .metal files, to understand what the chip could do. Then a first complete engine, in Rust, on a MacBook M5, with a Gemma model. It gave the right tokens. Against MLX, Apple's library, it got crushed on long prompts, by a factor close to ten :'(

I did not understand why. So I took MLX's routines one by one, compiled them as they were and measured them against mine, brick by brick. Attention first, because it was obvious. Yet there I found parity: I was looking in the wrong place :-/ Then the matrix multiplications, where the real gap was, even though I thought I had done what was needed.

Out of that came dozens of versions, in Rust then in C++. The last ones, in C++, called Metal directly and relied on part of MLX's routines, with my own kernels around them. They ran clearly faster than mlx-lm at prefill on a Gemma 4; at decode, it was a tie. That code is a bit old, it is not public and its measurements do not follow the series' protocol: no precise figure here.

Then I got fed up. The "intellectual" project of this series is to understand in detail. So I started over from scratch, with the series' emitter and nothing from MLX. That engine is the one this article measures.

Those months taught me what later served under Vulkan and under CUDA. A multiplication tile has to be measured on each machine: the one that wins for MLX does not necessarily win for me. From one chip to the next, the best one is not necessarily the same either. Attention has to go through the matrix units, otherwise long prompts crush it. That one has been my problem all along: like an idiot, a galactic-level black belt in impatience, I settled for small prompts, until I started working with my own engines. Finally, with unified memory, the GPU does not need the model copied over to it: it can read the file directly as the system mapped it into memory. On a 10 GB model, loading goes from several seconds to a few milliseconds. The disk read does not vanish, though: the system only reads a page of the file on first access, so the first pass pays for it, unless the file is already in memory.

The same tokens

The rule of the Vulkan article holds here too: speeds are only compared at equal tokens. For each model, two prompts, 48 tokens generated greedily, always taking the most likely one, compared with those of mlx-lm in bf16.

In f16, eleven sequences out of twelve are identical to the reference. In bf16, seven out of twelve; the other five depart from it along the way.

I noticed that where there is a divergence, most of the time the reference itself hesitates, and that is normal. I hinted at it in the CUDA article; it will be the subject of another article. For example, on the 27B, at the 16th token of a short prompt, the reference's two best candidates are separated only by the smallest gap bf16 can represent at that magnitude: a score of 27 for one, 26.875 for the other. In Herbert, both come out at 27. On a tie, the choice flips, and so the final text is not the same. Sometimes I used a 32-bit reference, to see what the model was really saying. In the other cases, the two candidates are already exactly tied in the reference. But the order of computations matters: with floating-point numbers, adding a and b then c does not always give exactly the same result as b and c then a. Herbert and MLX do not add in the same order. Between two tied candidates, that last bit is enough to tip the choice one way or the other.

Catching up on prefill

The new engine first ran with the Vulkan decode step ported as is. First campaign against four Mac engines: llama.cpp, mlx-lm, Rapid-MLX and BaseRT. llama.cpp led prefill on all six models. The next morning, same machine, same versions of the other engines, Herbert was ahead everywhere.

And yet the tokens produced did not change between the two.

Three things mattered:

1. The prefill matrix multiplications (the GEMMs) ran on the GPU's ordinary cores. I wrote them for its matrix units, the circuits that multiply small matrix blocks in one go, with simdgroup_matrix: the 0.8B's prefill went from 5,284 to 9,523 tokens per second. And since these GEMMs take 89% of the prefill's GPU time... can you picture it?

2. A question I had been asking myself for a long time: would f16 run faster than bf16 on this chip? Five probes later, the answer is no: as long as you accumulate in f32, nothing measurable changes. But the study found something else: matmul2d, Metal 4's matrix operation, beats my simdgroup_matrix GEMM by a third on the M3 Ultra. I did not expect that. Not by that much, anyway. Emitted by Herbert, it speeds up prefill by 13 to 34% depending on the model, without changing a single token. Yep, measuring matters :D

3. The engine no longer guesses which of the two GEMM forms, which tile (64 or 32 rows, 64, 32 or 16 columns), which block size to split the prompt, in which order to sort the experts: it measures at load time (did I ever tell you that you always have to measure? :D), on the machine, for each matrix shape of the model. It keeps the fastest. At that stage, for the dense models' GEMMs, it took 0.2 to 0.5 seconds. It has grown quite a bit since then, I come back to it at the end.

And then I took another blow to the head: the 35B-A3B is the one that moved the most. Its experts were computed one prompt row at a time: 523 tokens per second, four times behind llama.cpp. It is not fun every day, huh: just when you think you have a hold on some truth, reality tells you you're gonna need a bigger boat*. Grouped by expert then run as GEMMs, then in 16-row tiles, then measured at load time, they ended up at 2,356.

* "You're gonna need a bigger boat": Chief Brody (Roy Scheider) in Jaws (1975), on first seeing the shark. Reportedly ad-libbed on set. Said when a problem turns out much bigger than you thought.

That is the thread of the series: a compiler that knows the machine at the moment it emits. On the processor, under Vulkan or CUDA, it emitted for the machine it saw. Under Metal, no exception: it has to measure as well before choosing.

The night the engine chose wrong

Keeping the fastest without any other check is a trap. Naturally, I fell into it during the measurements for this article, right in the night before I wrote it. You cannot imagine how much I hate myself in those cases :D

Here is the thing: in the meantime, I had written kernels for the MacBook M5 that go through its new matrix units. On the M3 Ultra, they build without a problem and even run faster than the others there. So the engine picked them. Me, happy, obviously, huh.

Yeah, except that there they return wrong values: the way a matrix's elements are distributed across threads is not the same on the two chips. Rhaaaaaaa X-(

So, well, nothing serious, it was just the token comparison that showed it, in the middle of the night: zero identical tokens out of 48, for the bf16 engine as for the f16 engine. The speeds measured until then were those of wrong kernels. I threw them away... when I thought I had found a brilliant idea :-/

The fix: each candidate is run once against a kernel already in production before it is timed, on a full block and on a partial block, with the output zeroed beforehand. If it departs from it by more than 1/64 of the largest value, it is eliminated and the log says so. It took me two tries, of course, because the first check let a wrong shape through on the output projection of the 4B and the 9B. With the second, all of the M5's shapes are ruled out on the M3 Ultra. All of Herbert's measurements in this article were redone afterwards, at the published commit.

Here is what the 0.8B's load log shows for attention, on a 512-row block (raw lines, the last one cut at the ...). First the check, which rules out the M5 variant. Then the four timed variants: the M5's, marked with a star, was the fastest. mma is the one kept.

[autotune] sdpa 512 rows at 512: wrong against flash, never picked: nax (512 rows: max err 7.790e-2 of 9.082e-2)
[autotune] sdpa 512 rows at 512: flash 1.444 ms, mma 0.533 ms, mm2d 0.825 ms, nax* 0.461 ms -> mma
[metal] 688 kernels compiled in 0.11 s, GEMMs and blocks timed in 24.11 s, 137 dispatches per token, blocks of 2048 rows; ...

The attention variants in this log

• flash: the attention written at the start. It goes through the keys in blocks and computes the softmax on the fly, the FlashAttention method, but in scalar multiply-adds, one value at a time.
• mma: the same method on the classic matrix units (simdgroup_matrix), which multiply 8 × 8 blocks in one go.
• mm2d: the same again, with matmul2d, Metal 4's tensor operations.
• nax: the variant written for the M5's Neural Accelerators. Correct on the M5, wrong on the M3 Ultra.

The driver compiles the 688 kernels in 0.11 seconds. It is the measurements that take time: 24 seconds.

That costs one more pass per candidate at load time. I come back to it below.

Decode

In the first campaign, Herbert's decode was on par with llama.cpp, behind BaseRT: 238 tokens per second on the 0.8B, against 307. It was still the Vulkan step ported as is, one operation per kernel launch on the GPU (a dispatch): 263 launches per token on the 0.8B. I measured it first in the quantized version of the engine (still not published, but in production at our place, except that I am going to have to rework it entirely): the fixed cost of each launch weighed half the time of the projections there. The fusions that follow were born there, then I carried them over to the 16-bit engine.

I fused four modules: the norms into the projection that follows them, the SwiGLU into the gate_up projection, the whole DeltaNet recurrence step (the linear layers of Qwen3.5) into a single launch, and the RoPE into the attention.

A fusion means doing in a single kernel what two or three kernels did one after the other. Separate, each one writes its result to memory, the next one reads it back and each launch pays its fixed cost. Fused, the intermediate result stays in the GPU's registers: one launch instead of several, one write and one read fewer. Take the SwiGLU: without fusion, the gate_up projection writes its two outputs, then a second kernel reads them back to combine them; fused, the projection combines them before writing. On the 0.8B, the step went from 263 to 137 launches per token.

Each fusion gives the same bits as the kernels it replaces. And each one is an option the engine measures at load time: the fused norm wins on the 0.8B and the 2B, it loses from the 4B on, where each group reads back more data than weights.

These decode rates are measured in real time (wall-clock time), like those of llama.cpp and mlx-lm. Herbert's bench also gives the time spent on the GPU alone, 1 to 12% shorter at decode: I do not use it for comparisons. At prefill, the gap between the two stays under 0.5%. For mlx-lm, the 0.8B to 9B figures come from a benchmark where its 128 tokens follow a 2,048-token prompt. I measured them again from an empty context: they do not move, within 2%. The 27B and 35B-A3B figures come from that second measurement, where mlx-lm finally loaded the 27B. BaseRT did not run in this last campaign. In the previous one, on an earlier version of the engine, it did 299 tokens per second on the 0.8B, Herbert in f16 342. I did not have time to get to it while preparing this article for publication.

Under CUDA, the previous article had found a decode attention that fell behind on long contexts. On the Mac Studio, the lead holds at 32,768 tokens. It also holds after 81,920 tokens: 185 tokens per second on the 0.8B against 156 for mlx-lm and 154 for llama.cpp, 36.5 against 34 and 33 on the 9B. On the MacBook Air M1, that is not the case: after 81,920 tokens, Herbert decodes there at 15.3 tokens per second, llama.cpp at 21.5. I have not fixed it yet. Three test machines, tons of cases and still only 24 hours in a day...

f16

bf16 has been the series' format since the first article. Already under Vulkan, I had seen that it was not the ideal format everywhere. On the M1, bf16 has no fast path: its GPU computes it markedly more slowly than f16 (llama.cpp runs 18% faster there in f16, Herbert 40 to 50%). As for its processor's matrix unit, the AMX, it does not know it at all: I probed it. So Herbert has a second mode, built from the same code: bf16 weights converted to f16 at load time, round to nearest; everything else in f16, but accumulators still in f32. I also wanted to accumulate in f16, to gain even more time. The measurement said no: 1.5 to 4% on the matrix multiplications depending on the chip, nothing at all at decode, where memory is the limit. And much larger errors. I did not keep it. The GPU's f32-to-f16 conversion was checked on all 4.3 billion f32 values, on the M3 Ultra and on the M1. Against mlx-lm in float16, on the six Qwen models and two prompts, eleven sequences out of twelve are identical; the twelfth falls on an exact tie in the reference.

On the Mac Studio, f16 changes nothing: from 3% slower to 1% faster depending on the model. But on the MacBook Air M1, at a 512-token prefill, it speeds Herbert up by 40% on the 0.8B and 50% on the 2B. There, Herbert moves ahead of llama.cpp in f16 (1,484 tokens per second against 1,274 on the 0.8B, 657 against 564 on the 2B).

f16 is not as good as bf16 everywhere. It is more precise but covers a narrower range: a model trained in bf16 can contain values f16 cannot represent. On the Qwen models measured, the tokens say it works. I keep bf16 as the series' contract. f16 is an option for chips that have no bf16 or get nothing out of it. Really, I need to write an article on this subject.

Long prompts

I then ran the campaign up to 81,920 tokens. Reading the figures, I wrote: "Herbert is getting wrecked at prefill ... and my gut feeling here is that beyond a certain pp (prompt length) you have to use simdgroups and other hardware acceleration, on the M1, M3 Ultra and M5".

On the Mac Studio, at 81,920 tokens, the 0.8B did 1,128 tokens per second at prefill, llama.cpp 3,240 and mlx-lm 4,848. Up to 2,000 tokens, Herbert was ahead of everyone; at 4,000, mlx-lm and BaseRT moved ahead. Ouch.

The kernel-by-kernel profile, with Herbert's profiler that times each kernel on the GPU, showed the problem came from attention. On the 0.8B, its share of prefill went from 21% at 2,048 tokens to 89.5% at 61,440. The GEMMs, for their part, did not move.

As I said above, reading the code, the cause was simple: while the GEMMs had run on the matrix units from the start, the prefill attention had stayed in scalar multiply-adds. And here, the attention computation grows with the square of the prompt length. On a short prompt, it did not show.

And yet that is what I explain every week in the training courses I give. I was really furious with myself. I do not even remember why I wrote it that way. It was already the first problem I had run into with the C++ version, months earlier. Rewriting the engine, I did it again. That is a bit much, isn't it? What I do know is that I am going to have to go back over all my engines to check them again.

So I rewrote attention on the matrix units: the products of Q by K and of P by V in simdgroup_multiply_accumulate, 16-bit operands, f32 accumulators. The probabilities rounded to 16 bits at the same place as in the reference code. Q kept in registers. A matmul2d variant too, slower on the Mac Studio and the M1, level on the M5 at 2,048 tokens. As for the GEMMs, the engine measures the variants at load time and keeps the fastest.

On the 0.8B, attention's share of prefill drops from 21 to 6% at 2,048 tokens, from 89.5 to 60% at 61,440 tokens. It remains the biggest share on long prompts.

Here is an excerpt of what Herbert writes at startup for this attention, in f16, on the 0.8B (file sdpa_mma_h8kv2d256_q16k32_st_qr.metal, raw lines): four 8 × 8 accumulators in f32, then the product of a block of Q by a block of K, read from shared memory.

simdgroup_float8x8 m34 = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
simdgroup_float8x8 m35 = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
simdgroup_float8x8 m36 = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
simdgroup_float8x8 m37 = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
simdgroup_half8x8 v177;
simdgroup_load(v177, (threadgroup half *)s3 + 0u, 264, ulong2(0, 0), true);
simdgroup_multiply_accumulate(m34, v56, v177, m34);
simdgroup_half8x8 v179;
simdgroup_load(v179, (threadgroup half *)s3 + 2112u, 264, ulong2(0, 0), true);
simdgroup_multiply_accumulate(m35, v56, v179, m35);

The kernel's name is based on its shape: 8 heads, 2 K/V heads, dimension 256, tiles of 16 queries over 32 keys. It is one of the 688 kernels the engine emits when loading this model.

Here are the results of the measurements, in bf16, against mlx-lm and llama.cpp in bf16, two runs per measurement, engine order alternated:

At 81,920 tokens, the 0.8B goes from 1,128 to 4,827 tokens per second. The earlier measurement was in f16, this one in bf16; on the Mac Studio, both modes run at the same speed. Herbert is faster than llama.cpp at every length, by 18 to 49%. With mlx-lm, it is level at 61,440 tokens, within 5%, with the 2B 1% behind. The 4,848 tokens per second for mlx-lm quoted earlier, like the 3,240 for llama.cpp, come from the previous campaign; in this one, they do 4,784 and 3,232.

The two big models were measured at 512 and 8,192 tokens:

mlx-lm did not load the 27B in bf16, twice in a row ("GPU Timeout"); I do not know why yet. It loaded it a few hours later, for the decode measurement. In the first campaign, it had loaded it on the second try: 286 tokens per second at 512 tokens.

On the MacBook Air M1, in f16, the same attention multiplies prefill by four at 61,440 tokens: from 115 to 460 tokens per second, ahead of llama.cpp in f16 (401). This measurement was made on an earlier version, before the M5 kernels were added.

What I do not explain yet

The lead melts away with length. From 2,048 to 61,440 tokens, Herbert goes from ×1.21 to ×1.01 of mlx-lm on the 0.8B. Well, that is not totally crazy, huh: we are reaching the limits of the machine. But I still have a hypothesis: my attention remains slower than MLX's; it is the one that dominates on long prompts. I have not checked it kernel against kernel yet. 24-hour days and all that...

The MacBook Air M5. Its matrix units are the reason for the kernels that skewed the night of measurements. There, I measured Herbert ahead of mlx-lm on the 0.8B and the 2B at 2,048 tokens, level on the 4B, then at 0.81 times mlx-lm on the 4B at 8,192 tokens. Looking for why, I found the heat, the throttling I mentioned above. Cooled down, the 4B comes back level. I do not have a protocol that holds on this machine yet; its figures do not go in the tables. I will come back to it in another article, probably when I talk about the quantized versions of these engines.

The limits

One machine for the tables: a Mac Studio M3 Ultra, 96 GB, macOS 26.5. Six Qwen models, in bf16, one request at a time. Two runs per measurement. Long prompts were only measured up to the 9B. No quantization: that is another block of the series.

The mlx-lm and llama.cpp measurements of the last campaign date from its first pass, Herbert's from the pass redone after the fix, on the same machine, a few hours later. Rapid-MLX and BaseRT only ran in the previous campaigns, on an earlier version of the engine.

Gemma models also run on this engine, but their prompt still goes through token by token: their prefill is dozens of times behind mlx-lm. That is already no longer the case for the Q4 version of the engine, but I do not talk about it here. I am just going to do my best to update all my 16-bit engines and add support for several models to them.

These figures hold at the time this article is published, with these versions: llama.cpp fb4b2737, mlx-lm 0.32.0 (on mlx 0.32.3), Rapid-MLX 0.16.10, BaseRT 0.4.2. Engines evolve: some of these figures may no longer be true when you read this.

So, does it start up more slowly?

Yes. We compile Metal at startup, so the question comes up. But compilation is not the problem: the driver compiles the 0.8B's 688 kernels in 0.11 seconds. What takes time is measurement: timing each variant on each matrix shape, then checking it against a kernel in production. On the 0.8B, 20 to 35 seconds depending on the planned context length. On the 27B, more than two minutes.

For a command-line try, that is long. For a server that runs for days, it does not matter. In between, a cache is possible: the choices depend only on the chip, the driver, the engine version and the model's shape. You measure once, keep the result and measure again when one of the four changes. I have not written it yet.

It remains to be seen whether it is necessary and whether it is safe. On the M5, heat already skews the measurements from one minute to the next: a choice made cold and kept in a cache might no longer be the right one hot. I need time to think about it.

Reproducing at home

Two repositories are online, under the MIT license. The engine: github.com/xigh/herbert-metal-bf16, tag v0.1.0. bf16 is the default mode; f16 is enabled at build time.

git clone --branch v0.1.0 https://github.com/xigh/herbert-metal-bf16.git
cd herbert-metal-bf16 && cargo build --release --locked
cargo build --release --locked --features f16 --target-dir target/f16

You need an Apple Silicon Mac and the stable Rust toolchain, nothing else: Metal comes with the system. The model is a Hugging Face folder of Qwen3.5 (config.json, tokenizer.json and the bf16 weights in safetensors). To chat with it:

target/release/herbert-metal-chat ~/models/Qwen3.5-0.8B --no-system --no-think --ask 'What is the capital of France?'
target/f16/release/herbert-metal-chat ~/models/Qwen3.5-2B --no-system --no-think --ask 'Bonjour'

The default setting is the article's: everything is measured at load time, with the correctness check. To replay the old attention, the one that got wrecked, add --tune sdpa.prefill=1.

The campaign scripts, as they ran, the prediction written before the long prompts, the protocol and the raw data: github.com/xigh/herbert-metal-bench, tag v1-2026-11-01. If a figure looks wrong to you, you have what you need to check it.

What comes next

For the next article, I do not know yet exactly what I am going to talk to you about. We have just spent a month talking about CPUs, Vulkan, CUDA and now Metal, with really encouraging results, where Herbert holds its own against its buddies and often does better. On our side, we already have quantized versions running in production, but I have to revisit them, because the series of articles I have just written shed yet more light on details I thought I had understood. I can still tell you about plenty of engines, on Raspberry Pi, on old PCs without AVX, on Tenstorrent's Blackhole or on a RISC-V processor I designed on an FPGA.

But I am also making progress on the harness that goes with Herbert, which is also in production at our place and at a few of our customers and friends. And holy shit, it is really worth its weight in gold... I am working on the last details that will, I hope, allow a wide release. For me, that is now the priority. We have things to show and we are proud of them... so we will see how the week goes.

Thanks again for following me. It really means a lot to me. These articles are not easy for everyone to read, and yet I would like to make them accessible to many. The audience is limited; that is why, on my YouTube channel, on Instagram and Facebook, I have started doing a kind of popularization that seems much more accessible to me... When you get a chance, have a look, it is short. And tell me what you think!!!