On descend d'un cran : #004 le même compilateur sur les puces d'Apple

L'article précédent se terminait sur une question : est-ce que concevoir un compilateur pour moteur d'inférence LLM fonctionne sur les puces d'Apple ? Parce que dans ce cas précis, il n'est pas question d'écrire une sorte d'assembleur comme je l'avais fait pour les CPU, pour Vulkan avec SPIR-V ou pour les GPU NVIDIA avec PTX.

Ici, la seule option, c'est Metal, l'API d'Apple, sur des machines à mémoire unifiée (partagée entre le processeur et le GPU). Metal a son propre langage, le Metal Shading Language : un langage source, proche du C++, que le pilote compile lui-même. Apple ne documente ni le jeu d'instructions de ses GPU ni la forme intermédiaire que produit son compilateur.

L'info existe bien en partie, avec Asahi Linux, un projet incroyable de rétro-ingénierie : ses développeurs ont documenté une bonne partie du GPU d'Apple pour écrire des pilotes libres et faire tourner Linux sur les premières générations de Mac Apple Silicon. Mais rien d'utilisable pour nos besoins. Ici, Herbert n'écrit donc pas d'assembleur : il écrit du texte source au démarrage et le donne au pilote.

Et ce n'est pas tout. J'ai travaillé avec trois machines Apple. Pourtant, elles n'ont presque rien en commun.

En bref

• Sur un Mac Studio M3 Ultra, Herbert donne les mêmes tokens que mlx-lm, aux égalités près. Jusqu'à 8 192 tokens, il est devant mlx-lm et llama.cpp au prefill sur les six modèles, de 4 à 31 % sur mlx-lm (qui n'a pas chargé le 27B). Il est devant au décodage partout, de 9 à 34 % sur le plus rapide des deux ; l'avance tient après un prompt de 32 768 tokens.
• Aux premières campagnes de ce nouveau moteur Herbert, llama.cpp était devant au prefill sur les six modèles. Depuis, j'ai minutieusement adapté Herbert pour mesurer au chargement plusieurs formes de chaque noyau et garder la plus rapide. Pendant un temps, il choisissait des noyaux plus rapides... mais faux. Après un énième correctif, Herbert vérifie chaque candidat avant de le prendre.
• Évidemment, sur les prompts longs, Herbert s'est fait déboîter : à 81 920 tokens, 1 128 tokens par seconde contre 4 848 pour mlx-lm. L'attention du prefill était restée en calcul scalaire :-/ Réécrite sur les unités matricielles de la puce, elle donne 4 827, à égalité avec mlx-lm (4 784 dans la même campagne). L'avance fond avec la longueur : à 61 440 tokens, le 2B est 1 % derrière.

1. Un Mac Studio M3 Ultra, le gros de la bande : 96 Go de mémoire, 819 Go/s sur le papier pour la lire, de quoi tenir en bf16 des modèles qu'aucune carte grand public ne charge. Avec lui, la question est de ne rien gaspiller : un GPU de cette taille attend vite ses données. Bon, j'ai beaucoup à dire sur le contrôleur mémoire là, parce que la bande passante mémoire n'est pas celle qu'on annonce et j'ai été assez surpris de voir que le CPU n'y avait pas totalement accès. J'en ferai sûrement un post ou un article.

2. Un MacBook Air M1, mon plus vieil appareil : 8 Go, une soixantaine de Go/s mesurés, incapable, du coup, de tenir plus qu'un 2B. C'était le début d'Apple dans les processeurs pour ordinateurs et ça se voit : l'unité matricielle de son processeur ne connaît pas le format de nombre bf16 (le format de toute la série) et son GPU le calcule lentement : llama.cpp lui-même y va 18 % plus vite en f16 qu'en bf16. Je reviendrai sûrement plus en détail sur ce sujet dans un prochain post, voire un article.

3. Un MacBook Air M5, le dernier venu : 32 Go et des unités de calcul matricielles que les deux autres n'ont pas. Apple les appelle des « Neural Accelerators » et en met une dans chaque cœur du GPU ; on y accède par les opérations sur tenseurs de Metal 4, comme matmul2d. Sur cette puce, ça multiplie des matrices deux à trois fois plus vite que les unités matricielles classiques. Par contre, MacBook Air oblige : pas de ventilateur. Du coup, il va très vite pendant trente secondes, puis il chauffe et ne mesure plus la même chose : c'est le throttling. Passé une certaine température, le système baisse la fréquence de la puce pour la protéger : tout ralentit.

L'idée a été d'écrire un seul Herbert qui s'adapte de lui-même aux trois : au chargement, il regarde sur quelle puce il tourne, mesure ce qu'elle sait faire et s'y ajuste. C'est le cœur de mon travail, mais cette fois-ci, en Metal. Les noyaux, les tuiles, la taille des blocs sont choisis sur la machine, pas à l'avance ; le format des nombres, bf16 ou f16, se choisit à la construction, selon ce que la puce sait faire. C'est tout le sujet de cet article.

Le texte source qu'écrit Herbert est fait pour la forme du modèle et pour la puce : chaque dimension connue y figure en constante. Le pilote le compile sans « fast math », le mode où le compilateur s'autorise à réordonner les calculs et à prendre des raccourcis d'arrondi : les opérations de base (addition, multiplication, multiplication-addition fusionnée) sont arrondies comme le prévoit la norme IEEE. Pas de fichier .metal écrit à l'avance, pas de MLX, pas de Metal Performance Shaders, la bibliothèque de calcul d'Apple : le même émetteur que sous Vulkan, avec le même contrat numérique (poids, activations, cache K/V et états de la récurrence en bf16, accumulateurs en f32). Et Herbert n'arrondit qu'aux endroits où le code de référence de Hugging Face arrondit aussi.

La mémoire du Mac Studio permet aussi un choix de modèles plus large. Sur la 4080 et ses 16 Go, je m'arrêtais au 4B. Ici, six Qwen tiennent en bf16 : Qwen3.5 0.8B, 2B, 4B et 9B, Qwen3.8-27B (un dense de 55 Go) et Qwen3.6-35B-A3B, un modèle à experts (MoE) dont 3 milliards de paramètres seulement sont dits « actifs » et servent à chaque token. J'ai aussi porté Gemma 4, le modèle avec lequel tout a commencé (j'en parle juste après), mais ses chiffres viendront plus tard.

Des mois de travail

Ce moteur n'est pas le premier que j'écris pour les puces d'Apple. C'est même par elles que j'ai commencé à écrire des noyaux GPU, des mois avant Vulkan et CUDA. Et c'est le travail qui a été le plus difficile, si je mets à part celui que je mène sur Blackhole, la carte de Tenstorrent, dont je parlerai quand j'aurai fini.

J'ai d'abord écrit mes multiplications de matrices à la main, dans des fichiers .metal, pour comprendre ce que la puce savait faire. Puis un premier moteur complet, en Rust, sur un MacBook M5, avec un modèle Gemma. Il donnait les bons tokens. Face à MLX, la bibliothèque d'Apple, il se faisait écraser sur les prompts longs, d'un facteur proche de dix :'(

Je ne comprenais pas pourquoi. Alors j'ai pris les routines de MLX une par une, je les ai compilées telles quelles et je les ai mesurées contre les miennes, brique par brique. L'attention d'abord, parce que c'était une évidence. Pourtant, j'y ai trouvé la parité : je cherchais au mauvais endroit :-/ Puis les multiplications de matrices, où était le vrai écart, alors que je pensais avoir fait le nécessaire.

De là sont sorties des dizaines de versions, en Rust puis en C++. Les dernières, en C++, appelaient directement Metal et s'appuyaient sur une partie des routines de MLX, avec mes propres noyaux autour. Elles allaient nettement plus vite que mlx-lm au prefill sur un Gemma 4 ; au décodage, c'était l'égalité. Ce code date un peu, il n'est pas public et ses mesures ne suivent pas le protocole de la série : pas de chiffre précis ici.

Puis j'en ai eu marre. Le projet « intellectuel » de cette série, c'est de comprendre dans le détail. Alors j'ai tout repris de zéro, avec l'émetteur de la série et sans rien de MLX. C'est ce moteur-là que mesure cet article.

Ces mois m'ont appris ce qui a servi ensuite sous Vulkan et sous CUDA. Une tuile de multiplication se mesure sur chaque machine : celle qui gagne chez MLX ne gagne pas forcément chez moi. D'une puce à l'autre, la meilleure n'est d'ailleurs pas forcément la même. L'attention doit passer par les unités matricielles, sinon les prompts longs l'écrasent. Ce point-là a été mon problème tout au long de mon travail : comme un crétin, ceinture noire d'impatience niveau galactique, je me contentais de petits prompts, jusqu'à ce que je commence à bosser avec mes propres moteurs. Enfin, avec la mémoire unifiée, le GPU n'a pas besoin qu'on lui recopie le modèle : il peut lire directement le fichier tel que le système l'a projeté en mémoire. Sur un modèle de 10 Go, le chargement passe de plusieurs secondes à quelques millisecondes. La lecture du disque ne disparaît pas pour autant : le système ne lit une page du fichier qu'au premier accès, donc c'est le premier passage qui la paie, si le fichier n'est pas déjà en mémoire.

Les mêmes tokens

La règle de l'article sur Vulkan tient ici aussi : on ne compare des vitesses qu'à tokens égaux. Pour chaque modèle, deux prompts, 48 tokens générés en prenant toujours le plus probable (en « greedy », comme on dit), comparés à ceux de mlx-lm en bf16.

En f16, onze suites sur douze sont identiques à la référence. En bf16, sept sur douze ; les cinq autres s'en écartent en cours de route.

J'ai noté que là où il y a une divergence, le plus souvent la référence elle-même hésite et c'est normal. Je vous en parlais à demi-mot dans l'article sur CUDA ; ça sera l'objet d'un autre article. Par exemple, sur le 27B, au 16e token d'un prompt court, les deux meilleurs candidats de la référence ne sont séparés que par le plus petit écart que le bf16 sait représenter à cette hauteur : un score de 27 pour l'un, de 26,875 pour l'autre. Chez Herbert, les deux arrivent à 27. À égalité, le choix bascule et du coup, on n'a pas le même texte final. Parfois, j'ai utilisé une référence en 32 bits, pour voir ce que disait vraiment le modèle. Dans les autres cas, les deux candidats sont déjà à égalité exacte dans la référence. Mais l'ordre des calculs compte : avec des nombres à virgule flottante, additionner a et b puis c ne donne pas toujours exactement le même résultat que b et c puis a. Herbert et MLX n'additionnent pas dans le même ordre. Entre deux candidats à égalité, ce dernier bit suffit à faire pencher le choix d'un côté ou de l'autre.

Rattraper le prefill

Le nouveau moteur a d'abord tourné avec le pas de décodage de Vulkan porté tel quel. Première campagne contre quatre moteurs du Mac : llama.cpp, mlx-lm, Rapid-MLX et BaseRT. Llama.cpp menait le prefill sur les six modèles. Le lendemain matin, même machine, mêmes versions des autres moteurs, Herbert était devant partout.

Et pourtant les tokens produits n'ont pas changé entre les deux.

Trois choses ont compté :

1. Les multiplications de matrices du prefill (les GEMM) tournaient sur les cœurs ordinaires du GPU. Je les ai écrites pour ses unités matricielles, les circuits qui multiplient de petits blocs de matrice d'un coup, avec simdgroup_matrix : le prefill du 0.8B est passé de 5 284 à 9 523 tokens par seconde. Et comme ces GEMM occupent 89 % du temps GPU du prefill... vous imaginez ?

2. Une question que je me posais depuis longtemps : le f16 irait-il plus vite que le bf16 sur cette puce ? Cinq sondes plus tard, la réponse est non : tant qu'on accumule en f32, rien de mesurable ne change. Mais l'étude a trouvé autre chose : matmul2d, l'opération matricielle de Metal 4, bat ma GEMM simdgroup_matrix d'un tiers sur le M3 Ultra. Je ne m'y attendais pas. Pas à ce point en tout cas. Émise par Herbert, elle accélère le prefill de 13 à 34 % selon le modèle, sans changer un token. Ouais, la mesure compte :D

3. Le moteur ne devine plus laquelle des deux formes de GEMM, quelle tuile (64 ou 32 lignes, 64, 32 ou 16 colonnes), quelle taille de bloc pour découper le prompt, dans quel ordre trier les experts : il mesure au chargement (je vous ai déjà dit qu'il faut toujours mesurer ? :D), sur la machine, pour chaque forme de matrice du modèle. Il garde la plus rapide. À ce stade, pour les GEMM des denses, ça prenait 0,2 à 0,5 seconde. Ça a bien grossi depuis, j'y reviens à la fin.

Et puis j'ai encore pris un coup sur la tête : le 35B-A3B est celui qui a le plus bougé. Ses experts étaient calculés ligne par ligne du prompt : 523 tokens par seconde, quatre fois derrière llama.cpp. C'est pas drôle tous les jours, hein : quand tu crois tenir une vérité, le réel vient te montrer qui c'est Raoul*. Regroupés par expert puis passés en GEMM, puis en tuiles de 16 lignes, puis mesurés au chargement, ils ont fini à 2 356.

* « Montrer qui c'est Raoul » : réplique de Raoul Volfoni, joué par Bernard Blier, dans Les Tontons flingueurs (1963), dialogues de Michel Audiard. Remettre quelqu'un à sa place, sans ménagement.

C'est le fil de la série : un compilateur qui connaît la machine au moment où il émet. Sur le processeur, sous Vulkan ou CUDA, il émettait pour la machine qu'il voyait. Sous Metal, pas d'exception, il doit mesurer aussi avant de choisir.

La nuit où le moteur a choisi faux

Garder le plus rapide sans autre contrôle est un piège. Bien entendu, j'y suis tombé pendant les mesures de cet article, carrément la nuit qui précédait son écriture. Vous n'imaginez pas à quel point je me déteste dans ces cas-là :D

Voilà le truc : entre-temps, j'avais écrit pour le MacBook M5 des noyaux qui passent par ses nouvelles unités matricielles. Sur le M3 Ultra, ça se construit sans problème et ils y tournent même plus vite que les autres. Le moteur les a donc choisis. Moi, content, du coup, hein.

Ouais, sauf qu'ils y rendent des valeurs fausses : la façon dont les éléments d'une matrice sont répartis entre les fils n'est pas la même sur les deux puces. Rhaaaaaaa X-(

Alors bon, rien de grave hein, c'est juste la comparaison des tokens qui l'a montré, en pleine nuit : zéro token identique sur 48, pour le moteur bf16 comme pour le moteur f16. Les vitesses mesurées jusque-là étaient celles de noyaux faux. Je les ai jetées... alors que je pensais avoir trouvé l'idée géniale :-/

Le correctif : chaque candidat est exécuté une fois contre un noyau déjà en production avant d'être chronométré, sur un bloc entier et sur un bloc partiel, sortie remise à zéro avant. S'il s'en écarte de plus de 1/64 de la plus grande valeur, il est éliminé et le journal le dit. Il m'a fallu deux essais, bien entendu, parce que le premier contrôle laissait passer une forme fausse sur la projection de sortie du 4B et du 9B. Avec le second, toutes les formes du M5 sont écartées sur le M3 Ultra. Toutes les mesures d'Herbert de cet article ont été refaites après, au commit publié.

Voici ce que le journal de chargement du 0.8B affiche pour l'attention, sur un bloc de 512 lignes (lignes brutes, la dernière coupée aux ...). D'abord le contrôle, qui écarte la variante du M5. Puis les quatre variantes chronométrées : celle du M5, marquée d'une étoile, était la plus rapide. C'est mma qui est retenue.

[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; ...

Les variantes de l'attention dans ce journal

• flash : l'attention écrite au départ. Elle parcourt les clés par blocs et calcule le softmax au fil de l'eau, la méthode de FlashAttention, mais en multiplications-additions scalaires, une valeur à la fois.
• mma : la même méthode sur les unités matricielles classiques (simdgroup_matrix), qui multiplient des blocs de 8 × 8 d'un coup.
• mm2d : la même encore, avec matmul2d, les opérations sur tenseurs de Metal 4.
• nax : la variante écrite pour les Neural Accelerators du M5. Juste sur le M5, fausse sur le M3 Ultra.

Le pilote compile les 688 noyaux en 0,11 seconde. Ce sont les mesures qui prennent du temps : 24 secondes.

Ça coûte un passage de plus par candidat au chargement. J'y reviens plus bas.

Le décodage

À la première campagne, le décodage d'Herbert était au niveau de llama.cpp, derrière BaseRT : 238 tokens par seconde sur le 0.8B, contre 307. C'était encore le pas de Vulkan porté tel quel, une opération par lancement de noyau sur le GPU (un « dispatch », comme on dit) : 263 lancements par token sur le 0.8B. Je l'ai mesuré d'abord dans la version quantifiée du moteur (toujours pas publiée, mais en prod chez nous, sauf que je vais carrément devoir la reprendre) : le coût fixe de chaque lancement y pesait la moitié du temps des projections. Les fusions qui suivent y sont nées, puis je les ai reportées dans le moteur 16 bits.

J'ai fusionné quatre modules : les normes dans la projection qui les suit, la SwiGLU dans la projection gate_up, toute l'étape de la récurrence DeltaNet (les couches linéaires des Qwen3.5) en un seul lancement, le RoPE dans l'attention.

Une fusion, c'est faire en un seul noyau ce que deux ou trois noyaux faisaient à la suite. Séparés, chacun écrit son résultat en mémoire, le suivant le relit et chaque lancement paie son coût fixe. Fusionnés, le résultat intermédiaire reste dans les registres du GPU : un lancement au lieu de plusieurs, une écriture et une relecture en moins. Prenez la SwiGLU : sans fusion, la projection gate_up écrit ses deux sorties, puis un second noyau les relit pour les combiner ; fusionnée, la projection les combine avant d'écrire. Sur le 0.8B, le pas est passé de 263 à 137 lancements par token.

Chaque fusion donne les mêmes bits que les noyaux qu'elle remplace. Et chacune est une option que le moteur mesure au chargement : la norme fusionnée gagne sur le 0.8B et le 2B, elle perd dès le 4B, où chaque groupe relit plus de données que de poids.

Ces débits de décodage sont mesurés en temps réel (à l'horloge murale, comme on dit), comme ceux de llama.cpp et de mlx-lm. Le banc d'Herbert donne aussi le temps passé sur le GPU seul, de 1 à 12 % plus court au décodage : je ne m'en sers pas pour comparer. Au prefill, l'écart entre les deux reste sous 0,5 %. Pour mlx-lm, les chiffres du 0.8B au 9B viennent d'un banc où ses 128 tokens suivent un prompt de 2 048 tokens. Je les ai remesurés depuis un contexte vide : ils ne bougent pas, à 2 % près. Le 27B et le 35B-A3B viennent de cette seconde mesure, où mlx-lm a fini par charger le 27B. BaseRT n'a pas tourné dans cette dernière campagne. À la précédente, sur une version antérieure du moteur, il faisait 299 tokens par seconde sur le 0.8B, Herbert en f16 342. Je n'ai pas eu le temps de m'y mettre pour préparer la publication de cet article.

Sous CUDA, l'article précédent avait trouvé une attention de décodage qui décrochait sur les contextes longs. Sur le Mac Studio, l'avance tient à 32 768 tokens. Elle tient aussi après 81 920 tokens : 185 tokens par seconde sur le 0.8B contre 156 pour mlx-lm et 154 pour llama.cpp, 36,5 contre 34 et 33 sur le 9B. Sur le MacBook Air M1, ce n'est pas le cas : après 81 920 tokens, Herbert y décode à 15,3 tokens par seconde, llama.cpp à 21,5. Je ne l'ai pas encore corrigé. Trois machines de test, des tonnes de cas et toujours seulement 24 heures par jour...

Le f16

Le bf16 est le format de la série depuis le premier article. Déjà sous Vulkan, j'avais vu que ce n'était pas le format idéal partout. Sur le M1, le bf16 n'a pas de chemin rapide : son GPU le calcule nettement plus lentement que le f16 (llama.cpp y va 18 % plus vite en f16, Herbert 40 à 50 %). Quant à l'unité matricielle de son processeur, l'AMX, elle ne le connaît pas du tout : je l'ai sondée. Herbert a donc un second mode, construit depuis le même code : les poids bf16 convertis en f16 au chargement, au plus proche ; tout le reste en f16, mais accumulateurs toujours en f32. J'ai aussi voulu accumuler en f16, pour gagner encore du temps. La mesure a dit non : de 1,5 à 4 % sur les multiplications de matrices selon la puce, rien du tout au décodage, où c'est la mémoire qui limite. Et des erreurs bien plus grandes. Je ne l'ai pas gardé. La conversion de f32 vers f16 du GPU a été vérifiée sur les 4,3 milliards de valeurs f32, sur le M3 Ultra et sur le M1. Contre mlx-lm en float16, sur les six Qwen et deux prompts, onze suites sur douze sont identiques ; la douzième tombe sur une égalité exacte dans la référence.

Sur le Mac Studio, le f16 ne change rien : de 3 % plus lent à 1 % plus rapide selon le modèle. Mais sur le MacBook Air M1, au prefill de 512 tokens, il accélère Herbert de 40 % sur le 0.8B et de 50 % sur le 2B. Herbert y passe devant llama.cpp en f16 (1 484 tokens par seconde contre 1 274 sur le 0.8B, 657 contre 564 sur le 2B).

Le f16 ne vaut pas le bf16 partout. Il est plus précis mais couvre une plage plus étroite : un modèle entraîné en bf16 peut contenir des valeurs que le f16 ne sait pas représenter. Sur les Qwen mesurés, les tokens disent que ça passe. Je garde le bf16 comme contrat de la série. Le f16 est une option pour les puces qui n'ont pas de bf16 ou qui n'en tirent pas parti. Vraiment, il faut que j'écrive un article sur ce sujet.

Les prompts longs

J'ai ensuite lancé la campagne jusqu'à 81 920 tokens. En lisant les chiffres, j'ai écrit : « Herbert se fait déboîter en prefill ... et là mon intuition, c'est qu'au-delà d'un certain pp (longueur de prompt) il faut utiliser les simdgroups et autres accélérations matérielles, sur le M1, M3 Ultra et M5 ».

Sur le Mac Studio, à 81 920 tokens, le 0.8B faisait 1 128 tokens par seconde au prefill, llama.cpp 3 240 et mlx-lm 4 848. Jusqu'à 2 000 tokens, Herbert était devant tout le monde ; à 4 000, mlx-lm et BaseRT passaient devant. Ouch.

Le profil noyau par noyau, avec le profileur d'Herbert qui chronomètre chaque noyau sur le GPU, a montré que le problème venait de l'attention. Sur le 0.8B, sa part du prefill passait de 21 % à 2 048 tokens à 89,5 % à 61 440. Les GEMM, elles, ne bougeaient pas.

Je vous l'ai dit plus haut, en relisant le code, la cause était simple : alors que les GEMM tournaient sur les unités matricielles depuis le début, l'attention du prefill était restée en multiplications-additions scalaires. Et ici, le calcul de l'attention croît avec le carré de la longueur du prompt. Sur un prompt court, ça ne se voyait pas.

Et pourtant, c'est ce que j'explique chaque semaine pendant les formations que je donne. Vraiment j'étais en pétard contre moi. Je ne me souviens même pas pourquoi j'ai écrit ça comme ça. C'était déjà le premier problème que j'avais rencontré avec la version C++, des mois plus tôt. En réécrivant le moteur, je l'ai refait encore. Ça fait beaucoup, là, non ? Ce que je sais, c'est que je vais devoir repasser sur tous mes moteurs pour vérifier à nouveau.

J'ai donc réécrit l'attention sur les unités matricielles : les produits de Q par K et de P par V en simdgroup_multiply_accumulate, opérandes 16 bits, accumulateurs f32. Les probabilités arrondies à 16 bits au même endroit que dans le code de référence. Q gardé en registres. Une variante en matmul2d aussi, plus lente sur le Mac Studio et le M1, à égalité sur le M5 à 2 048 tokens. Comme pour les GEMM, le moteur mesure les variantes au chargement et garde la plus rapide.

Sur le 0.8B, la part de l'attention dans le prefill tombe de 21 à 6 % à 2 048 tokens, de 89,5 à 60 % à 61 440 tokens. Elle reste la plus grosse part sur les prompts longs.

Voici un extrait de ce qu'Herbert écrit au démarrage pour cette attention, en f16, sur le 0.8B (fichier sdpa_mma_h8kv2d256_q16k32_st_qr.metal, lignes brutes) : quatre accumulateurs 8 × 8 en f32, puis le produit d'un bloc de Q par un bloc de K, lu en mémoire partagée.

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);

Le nom du noyau est basé sur sa forme : 8 têtes, 2 têtes K/V, dimension 256, des tuiles de 16 requêtes sur 32 clés. Il fait partie des 688 noyaux que le moteur émet au chargement de ce modèle.

Voici le résultat des mesures, en bf16, contre mlx-lm et llama.cpp en bf16, deux runs par mesure, ordre des moteurs alterné :

À 81 920 tokens, le 0.8B passe de 1 128 à 4 827 tokens par seconde. La mesure d'avant était en f16, celle-ci en bf16 ; sur le Mac Studio, les deux modes vont à la même vitesse. Herbert va plus vite que llama.cpp à toutes les longueurs, de 18 à 49 %. Avec mlx-lm, c'est l'égalité à 61 440 tokens, à 5 % près, avec le 2B 1 % derrière. Les 4 848 tokens par seconde de mlx-lm cités plus haut, comme les 3 240 de llama.cpp, viennent de la campagne précédente ; dans celle-ci, ils font 4 784 et 3 232.

Les deux gros modèles ont été mesurés à 512 et 8 192 tokens :

Mlx-lm n'a pas chargé le 27B en bf16, deux fois de suite (« GPU Timeout ») ; je ne sais pas encore pourquoi. Il l'a chargé quelques heures plus tard, pour la mesure du décodage. À la première campagne, il l'avait chargé au second essai : 286 tokens par seconde à 512 tokens.

Sur le MacBook Air M1, en f16, la même attention multiplie le prefill par quatre à 61 440 tokens : de 115 à 460 tokens par seconde, devant llama.cpp en f16 (401). Cette mesure a été faite sur une version antérieure, avant l'ajout des noyaux du M5.

Ce que je n'explique pas encore

L'avance fond avec la longueur. De 2 048 à 61 440 tokens, Herbert passe de ×1,21 à ×1,01 de mlx-lm sur le 0.8B. Bon, c'est pas totalement déconnant, hein : on atteint les limites de la machine. Mais j'ai quand même une hypothèse : mon attention reste moins rapide que celle de MLX ; c'est elle qui domine sur les prompts longs. Je ne l'ai pas encore vérifiée noyau contre noyau. Les journées de 24 heures, toussa...

Le MacBook Air M5. Ses unités matricielles sont la raison des noyaux qui ont faussé la nuit des mesures. J'y ai mesuré Herbert devant mlx-lm sur le 0.8B et le 2B à 2 048 tokens, à égalité sur le 4B, puis à 0,81 fois mlx-lm sur le 4B à 8 192 tokens. En cherchant pourquoi, j'ai trouvé la chaleur, le throttling dont je parlais plus haut. Refroidi, le 4B revient à égalité. Je n'ai pas encore de protocole qui tienne sur cette machine ; ses chiffres ne vont pas dans les tableaux. J'y reviendrai dans un autre article, sûrement quand je parlerai des versions quantifiées de ces moteurs.

Les limites

Une machine pour les tableaux : un Mac Studio M3 Ultra, 96 Go, macOS 26.5. Six Qwen, en bf16, une requête à la fois. Deux runs par mesure. Les prompts longs n'ont été mesurés que jusqu'au 9B. Pas de quantification : c'est un autre bloc de la série.

Les mesures de mlx-lm et de llama.cpp de la dernière campagne datent de son premier passage, celles d'Herbert du passage refait après le correctif, sur la même machine, quelques heures plus tard. Rapid-MLX et BaseRT n'ont tourné que dans les campagnes précédentes, sur une version antérieure du moteur.

Les modèles Gemma tournent aussi sur ce moteur, mais leur prompt passe encore token par token : leur prefill est des dizaines de fois derrière mlx-lm. Ça n'est déjà plus le cas pour la version Q4 du moteur, mais je n'en parle pas ici. Je vais juste faire au mieux pour mettre à jour tous mes moteurs 16 bits et leur ajouter le support de plusieurs modèles.

Ces chiffres valent à la publication de cet article, avec ces versions : llama.cpp fb4b2737, mlx-lm 0.32.0 (sur mlx 0.32.3), Rapid-MLX 0.16.10, BaseRT 0.4.2. Les moteurs évoluent : certains de ces chiffres ne seront peut-être plus vrais quand vous lirez ceci.

Du coup, ça démarre moins vite ?

Oui. On compile du Metal au démarrage, alors la question se pose. Mais la compilation n'est pas le problème : le pilote compile les 688 noyaux du 0.8B en 0,11 seconde. Ce qui prend du temps, c'est la mesure : chronométrer chaque variante sur chaque forme de matrice, puis la vérifier contre un noyau en production. Sur le 0.8B, de 20 à 35 secondes selon la longueur de contexte prévue. Sur le 27B, plus de deux minutes.

Pour un essai en ligne de commande, c'est long. Pour un serveur qui tourne des jours, ça ne compte pas. Entre les deux, un cache est possible : les choix ne dépendent que de la puce, du pilote, de la version du moteur et de la forme du modèle. On mesure une fois, on garde le résultat et on remesure quand l'un des quatre change. Je ne l'ai pas encore écrit.

Reste à savoir s'il est nécessaire et s'il est sûr. Sur le M5, la chaleur fausse déjà les mesures d'une minute à l'autre : un choix fait à froid et gardé en cache ne serait peut-être plus le bon à chaud. J'ai besoin de temps pour y réfléchir.

Reproduire chez soi

Deux dépôts sont en ligne, sous licence MIT. Le moteur : github.com/xigh/herbert-metal-bf16, tag v0.1.0. Le bf16 est le mode par défaut ; le f16 s'active à la construction.

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

Il faut un Mac Apple Silicon et la chaîne Rust stable, rien d'autre : Metal est fourni par le système. Le modèle est un dossier Hugging Face de Qwen3.5 (config.json, tokenizer.json et les poids bf16 en safetensors). Pour discuter avec lui :

target/release/herbert-metal-chat ~/models/Qwen3.5-0.8B --no-system --no-think --ask 'Quelle est la capitale de la France ?'
target/f16/release/herbert-metal-chat ~/models/Qwen3.5-2B --no-system --no-think --ask 'Bonjour'

Le réglage par défaut est celui de l'article : tout est mesuré au chargement, avec le contrôle de justesse. Pour rejouer l'ancienne attention, celle qui se faisait déboîter, ajoutez --tune sdpa.prefill=1.

Les scripts de la campagne, tels qu'ils ont tourné, la prédiction écrite avant les prompts longs, le protocole et les données brutes : github.com/xigh/herbert-metal-bench, tag v1-2026-11-01. Si un chiffre vous semble faux, vous avez de quoi le vérifier.

La suite

Pour le prochain article, je ne sais pas encore précisément ce dont je vais vous parler. Là, on vient de passer un mois à parler de CPU, de Vulkan, de CUDA et maintenant de Metal, avec des résultats vraiment encourageants, où Herbert fait jeu égal et souvent mieux que ses copains. De notre côté on a déjà des versions quantifiées qui tournent en prod, mais que je dois revoir tant la série d'articles que je viens d'écrire m'a encore éclairé sur des détails que je pensais avoir pourtant compris. Je peux encore vous parler de plein de moteurs, sur Raspberry Pi, sur des vieux PC sans AVX, sur Blackhole Tenstorrent ou bien sur un processeur RISC-V que j'ai conçu en FPGA.

Mais j'avance aussi sur le harnais qui accompagne Herbert qui est aussi en production chez nous et chez quelques-uns de nos clients et amis. Et putain, il vaut vraiment son pesant de cacahuètes... Je travaille sur les derniers détails qui vont permettre, je l'espère, une diffusion massive. Pour moi, c'est maintenant la priorité. On a des choses à montrer et on en est fiers... alors on verra comment se passe la semaine.

Merci encore de me suivre. Ça me touche vraiment. Ces articles ne sont pas simples à lire par tout le monde et pourtant je voudrais les rendre accessibles à beaucoup. L'audience est limitée ; c'est pour ça que sur ma chaîne YouTube, sur Instagram et Facebook, j'ai commencé à faire une sorte de vulgarisation qui me semble bien plus accessible... À l'occasion, jetez-y un œil, c'est court. Et dites-moi ce que vous en pensez !!!