On descend d'un cran : #003 ce que CUDA a et que Vulkan n'a pas

L'article précédent se terminait sur une question : ce que CUDA pourrait nous apprendre qu'on ne sait pas faire en Vulkan. J'avais écrit qu'il manquait à Vulkan deux instructions, sans les nommer. Cinq des sept relectures de l'article l'ont relevé. Voici la réponse : il n'en manque pas deux, il en manque une.

Pour les mesurer, il fallait un Herbert qui parle CUDA. Il écrit ses noyaux en PTX, le langage intermédiaire de NVIDIA, au démarrage, pour la forme du modèle et pour la carte. Il charge le pilote à l'exécution. Le pilote contient son propre assembleur, l'équivalent de ptxas du kit CUDA : c'est lui qui traduit ce PTX en code machine pour la carte, choisit les registres, puis garde le résultat en cache. ptxas lui-même ne sert qu'à vérifier les modules hors ligne. Pas de nvcc, pas de NVRTC, pas de cuBLAS : le même principe que sur le processeur et sous Vulkan, avec le même contrat numérique (poids, activations et cache en bf16, accumulateurs en f32).

En bref

• Ce que CUDA a et que Vulkan n'a pas se réduit, au code machine, à une instruction : cp.async. Retirée du moteur CUDA, elle coûte 5 à 14 % au prefill. L'autre, ldmatrix, le pilote Vulkan l'émet déjà de lui-même. Au décodage, aucune des deux ne compte.
• Sur une RTX 4080, une requête à la fois, Herbert sous CUDA décode plus vite que TensorRT-LLM, le moteur de NVIDIA, sur les trois modèles mesurés (de 2 à 7 %). Il va aussi plus vite que vLLM, SGLang et llama.cpp. Au décodage, l'avance tient après un prompt de 40 000 tokens. Au prefill, c'est l'égalité à quelques pour cent sur les prompts longs, dans les deux sens. Sur le 2B, Herbert n'est derrière que faute d'un réglage mesuré pour ce modèle.
• J'avais prédit que je me ferais « déboîter » sur les prompts très longs. C'était vrai : à 60 000 tokens, Herbert décodait à la moitié de la vitesse de vLLM. J'en ai trouvé la cause et je l'ai corrigée. Il est maintenant devant au décodage à cette longueur, à égalité au prefill.

Deux instructions

Sur un GPU NVIDIA, une multiplication de matrices bien écrite fait deux choses en même temps : elle calcule sur une tuile pendant qu'elle charge la suivante. Deux instructions du PTX servent à ça.

cp.async copie une tuile de la mémoire de la carte vers la mémoire partagée du bloc sans passer par les registres ni attendre : le calcul continue pendant la copie. ldmatrix charge en une instruction, depuis la mémoire partagée, les fragments de matrice dont les tensor cores ont besoin, dans la disposition qu'ils attendent.

Voici ce qu'Herbert écrit au démarrage pour le GEMM de la projection gate_up du 0.8B (fichier gemm_1024x7168_bf16_64x128x32_w2x2_s2.ptx, lignes brutes, coupées aux ...) : la copie asynchrone de la tuile suivante, puis le calcul sur la tuile courante, ldmatrix puis mma sur les tensor cores.

	cp.async.wait_group 0;
	bar.sync 0;
	cp.async.cg.shared.global [%r83], [%rd9], 16, %r16;
	cp.async.cg.shared.global [%r84], [%rd13], 16, %r25;
	...
	cp.async.commit_group;
	ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r91, %r92, %r93, %r94}, [%r89];
	ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r99, %r100, %r101, %r102}, [%r90];
	mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%f0, %f1, %f2, %f3}, {%r91, %r92, %r93, %r94}, {%r99, %r100}, {%f0, %f1, %f2, %f3};
	...

Le reste du noyau est dans le README du moteur.

Sous Vulkan, SPIR-V n'a ni l'une ni l'autre. La copie asynchrone existe dans SPIR-V, mais elle repose sur des types réservés aux noyaux OpenCL, interdits dans un shader Vulkan. ldmatrix n'a pas d'équivalent explicite : les matrices coopératives laissent le pilote choisir comment charger. C'est ce que j'avais en tête en écrivant l'article précédent.

Le code machine m'a donné tort à moitié. Je ne m'y attendais pas. J'ai lu le code SASS qu'exécute vraiment la carte pour les noyaux Vulkan d'Herbert, dans le cache disque du pilote. Le GEMM coopératif contient bien des LDSM : le pilote émet ldmatrix de lui-même. Il ne contient aucun LDGSTS, l'instruction de cp.async : chaque tuile passe par les registres. Au niveau du code machine, ce que CUDA a et que Vulkan n'a pas, c'est une instruction.

Ce qu'elle vaut

Restait à savoir combien elle vaut. Je l'ai mesuré à l'envers : dans le moteur CUDA complet, j'ai retiré un réglage à la fois, avec le même binaire, sur la 4080, en alternant les variantes sur cinq tours. Prefill de 128 à 4 000 tokens, sur le 0.8B et le 4B. Les tokens produits sont identiques dans toutes les variantes : seule la vitesse change.

cp.async pèse le plus. C'est justement celle qui manque. Au décodage, aucun réglage du GEMM ne compte : seul le graphe CUDA, qui enregistre le pas d'un token une fois pour toutes, pèse 4,7 % sur le 0.8B et 1,1 % sur le 4B. Vulkan a l'équivalent du graphe, avec ses command buffers enregistrés.

Une précision : cette mesure a tourné le 18 octobre, carte en P2, l'état où le pilote la garde sous CUDA (j'y reviens plus bas). Ce sont des écarts entre variantes du même moteur dans le même état, pas des vitesses à comparer aux autres moteurs.

Il y a un autre levier propre à CUDA, que je n'ai pas pris : accumuler en f16. L'article précédent le disait déjà, le format seul ne change presque rien. Sous CUDA aussi, Herbert en f16 fait exactement la vitesse du bf16, à 0,5 % près. Le seul gain vient d'accumuler les produits en f16 au lieu du f32 : 6 à 21 % au prefill. Mais c'est un autre contrat numérique : sur un GEMM isolé, 14 % des sorties changent. Je ne l'ai pas pris.

Mes prédictions, puis les mesures

Avant la campagne, le 18 octobre au matin, j'ai dicté mes prédictions. Elles sont publiées telles quelles avec les données, dans un fichier daté. Elles ne sont pas aveugles : je connaissais les premières mesures d'Herbert sous CUDA contre une version plus ancienne de vLLM. Mais je ne connaissais ni les nouvelles versions des moteurs, ni leurs meilleurs réglages. Je me voyais derrière eux, partout. L'argument était simple : ces moteurs existent depuis des années, écrits par des ingénieurs extrêmement compétents. Je ne voyais aucune chance que mon code aille au moins aussi vite.

Le banc : une RTX 4080, pilote 595.99.02, horloges verrouillées (cœur à 2 805 MHz, mémoire à 11 201 MHz), Qwen3.5-0.8B et 4B en bf16, une requête à la fois. En face : llama.cpp b11718, vLLM 0.31.0, SGLang 0.5.22 et transformers 5.20.1, chacun à son meilleur réglage, cherché avant la mesure. Tous les moteurs CUDA en P0. Deux instruments, comme dans l'article précédent : le banc de chaque moteur et un client HTTP commun.

À ma grande surprise, c'est l'inverse, presque partout.

Le client HTTP commun dit la même chose, avec un ordre parfois différent. Au prefill de 4 096 tokens sur le 4B, SGLang, Herbert et vLLM tiennent dans 5 % les uns des autres.

Je ne suis pas du genre à faire le kéké : il faut dire d'où vient une bonne partie de l'avance. Le moteur nu, sans le choix automatique des tuiles du GEMM, est de 15 à 28 % derrière vLLM et SGLang au prefill du 4B. Les tableaux donnent le moteur avec ce choix, qui est son réglage par défaut depuis le 19 octobre, dans la version publiée.

Face au moteur de NVIDIA

Il manquait un moteur à ce banc : TensorRT-LLM, celui que NVIDIA développe pour ses propres cartes. Je l'ai ajouté le 24 octobre au soir, pour cet article, avec un troisième modèle, le 2B. Même carte, mêmes horloges verrouillées, tout en bf16. Herbert à ses réglages par défaut : le moteur mesuré émet exactement le même PTX que la version publiée.

J'ai mesuré le prefill de 128 à 40 000 tokens, puis le décodage après des prompts de 759 à 40 000 tokens. La mesure venait de commencer quand j'ai dicté une prédiction de plus : « jeu égal en décode, voire un peu mieux, mais moins bien avec plus de pp ... et un peu moins bien en prefill, tout le temps ... c'est le code de Nvidia, même si j'imagine qu'ils s'occupent plus des H100, B200 etc ... je pense qu'ils savent mieux utiliser la puissance de leur hardware que moi ».

Ces chiffres ne se lisent pas tout à fait comme ceux de la partie précédente. TensorRT-LLM, vLLM et SGLang ont tourné chacun en un passage, avec leurs propres répétitions, après les tours alternés d'Herbert et de llama.cpp. Leurs prompts sont des tokens tirés au hasard, ceux d'Herbert un vrai texte. Pour TensorRT-LLM et vLLM, le prefill est la latence d'une requête entière qui génère un token, ordonnancement compris : sur 128 tokens, ce coût fixe pèse lourd : le ×3 du 0.8B en dit surtout ça. De 2 000 à 40 000 tokens, quand le calcul domine, les écarts sont de quelques pour cent, dans les deux sens. Sur le 2B, TensorRT-LLM, SGLang et vLLM gardent l'avantage. Sur ces prompts longs, la carte atteint sa limite de 320 W et son horloge descend par moments jusqu'à 2 625 MHz ; vLLM découpe le prompt en blocs de 8 192 tokens ; TensorRT-LLM doit faire de même sur le 4B à 40 000 tokens.

Au décodage, en revanche, la mesure ne laisse pas de doute. Là aussi, j'ai été surpris : sur une carte NVIDIA, sans bibliothèque NVIDIA, Herbert va plus vite que le moteur de NVIDIA.

Toute proportion gardée : les moteurs auxquels je me compare sont faits pour servir beaucoup d'utilisateurs à la fois ; ce banc ne mesure qu'une requête à la fois. Je n'ai pas mesuré les cartes de centre de données sur lesquelles travaillent leurs équipes. Je compte le faire avec RunPod et je pense que l'équilibre y sera restauré. Mais pour moi, c'est un moment important. Je ne pensais pas arriver aussi vite à un tel résultat.

Restait le 2B, le seul modèle où Herbert est derrière au prefill. La raison est dans le code : le choix automatique des tuiles du GEMM n'a été mesuré que pour les largeurs du 0.8B et du 4B. Le 2B, qui n'en a pas, retombe sur un réglage de base, avec des blocs plus petits. Avec la tuile du 0.8B et du 4B et des blocs de 2 048 lignes, fixées par les options de la version publiée, en alternant les variantes sur cinq tours, son prefill passe de 25 684 à 29 056 tokens par seconde à 2 000 tokens, puis de 24 524 à 27 706 à 10 000 : 5 et 6 % devant TensorRT-LLM. Les tokens produits sont les mêmes. Ce réglage n'est pas encore le défaut ; il le deviendra dans une prochaine version. Pour le rejouer :

--tune gemm.auto=0 --tune gemm.tile=64x128w2x2 --tune gemm.bk=32 --tune gemm.stages=2 --tune prefill.block=2048

Le prefill long

Le 19 octobre, avant de mesurer des prompts de 10 000 à 60 000 tokens, j'ai dicté une prédiction de plus : « Je crois que je vais me faire déboîter pour les préfils longs. Et le décode aussi d'ailleurs. »

J'avais raison. À 60 000 tokens sur le 0.8B, Herbert faisait 0,84 fois SGLang et 0,87 fois vLLM au prefill. Au décodage qui suit un prompt de cette longueur, c'était pire : 0,49 fois vLLM. De 10 000 à 60 000 tokens, Herbert perdait 55 % de sa vitesse de décodage, vLLM et llama.cpp environ 20 %.

Le profil a désigné un seul noyau, l'attention du décodage. Entre 10 000 et 60 000 tokens, il passait de 22 à 62 % du temps de chaque token ; le reste ne bougeait pas. En lisant le code, la cause était simple : ce noyau découpait le travail en 64 blocs, quelle que soit la longueur du contexte, pour une carte qui a 76 multiprocesseurs. Chaque bloc parcourait seul sa part des positions. Et les quatre têtes qui partagent le même K et le même V, en attention groupée, le relisaient chacune. Le noyau tenait 149 Go/s, le cinquième de ce que la mémoire de la carte peut donner. Je me suis détesté d'avoir laissé passer ça : ne lire chaque donnée qu'une fois, c'est pourtant tout le cœur de ma démarche.

La correction a pris une nuit. Un bloc par tête K/V et par tranche, K et V lus une seule fois pour les quatre têtes, 64 tranches au lieu de 8. Au prefill, la copie de la tuile suivante part pendant le calcul de la précédente. Sur le 0.8B à 60 000 tokens, le décodage passe de 133 à 278 tokens par seconde, le prefill gagne 17 %.

Puis la mesure contre les autres, dans la même session, en alternance, cinq tours, à 60 000 tokens :

Au décodage, Herbert est maintenant devant à cette longueur. Au prefill, c'est l'égalité, à 4 % près, avec llama.cpp et vLLM encore devant sur le 4B, SGLang sur le 0.8B. Ces corrections changent l'ordre de certaines sommes : sur un prompt du 0.8B, la suite de tokens s'écarte maintenant de la référence au 172e token, où elle était identique ; sur le 4B, une autre devient identique. Ces écarts tombent sur des quasi-égalités entre deux tokens, comme dans l'article précédent. La correction est devenue le réglage par défaut le 21 octobre. L'ancienne attention reste disponible par trois options, pour rejouer les mesures d'avant.

L'horloge mémoire

L'article précédent laissait une question ouverte : sous CUDA, la mémoire de la 4080 tournait à 10 801 MHz au lieu de 11 201. Je ne l'expliquais pas.

C'est voulu par le pilote. Dès qu'un programme ouvre un contexte CUDA, il garde la carte dans l'état P2, un cran sous son maximum, sous Linux comme sous Windows. La variable d'environnement qu'on trouve partout pour l'éviter sert en fait à l'inverse : elle coupe le boost. Chez moi, elle n'a rien changé. Ce qui marche, c'est un profil d'application que NVIDIA livre avec ses pilotes, CudaNoStablePerfLimit. Avec lui, la carte reste en P0 et le décodage du 4B gagne 3,6 %, le rapport des deux horloges. Toute la campagne de cet article a tourné en P0, par souci d'équité pour tous les moteurs CUDA, pas seulement pour Herbert. Sans le profil, les moteurs mesurés le 18 octobre (Herbert, vLLM, SGLang, llama.cpp et transformers) perdent tous de 2,7 à 3,7 % au décodage. TensorRT-LLM, ajouté ensuite, n'a tourné qu'en P0 : le relevé des horloges le confirme seconde par seconde pendant ses mesures, mais je ne l'ai pas mesuré en P2.

Ce que je n'explique pas encore

Le prefill au-delà de 20 000 tokens : l'attention du prefill y prend 62 % du temps, avec un seul bloc de quatre warps par multiprocesseur, 8 % de ce que la carte pourrait faire tourner en même temps. Je n'ai pas encore trouvé comment la remplir sans changer les résultats.

Le décodage d'Herbert sous Vulkan sur le 0.8B en contexte long, qui décrochait déjà dans l'article précédent : la cause trouvée sous CUDA y est probablement aussi. Je ne l'ai pas encore mesuré.

Les limites

Une carte, la 4080. Trois modèles denses, en bf16. Une requête à la fois. Pas de B200, pas de H100, pas de FP8, pas de quantification. vLLM et SGLang sont faits pour servir beaucoup de requêtes à la fois : ce banc ne mesure pas ce qu'ils font de mieux.

Ces chiffres valent du 18 au 25 octobre 2026, avec ces versions et ce pilote (TensorRT-LLM 1.4.0rc1). Les moteurs évoluent : certains de ces chiffres ne seront peut-être plus vrais quand vous lirez ceci.

Reproduire chez soi

Deux dépôts sont en ligne, sous licence MIT. Le moteur : github.com/xigh/herbert-cuda-bf16, tag v0.1.0. Il se construit sans le SDK CUDA : il ne lui faut que le pilote NVIDIA.

git clone --branch v0.1.0 https://github.com/xigh/herbert-cuda-bf16.git
cd herbert-cuda-bf16 && cargo build --release --locked
target/release/herbert-cuda-chat ~/models/Qwen3.5-4B --no-think --ask 'Quelle est la capitale de la France ?'

Le réglage par défaut est celui de l'attention corrigée. Pour rejouer les mesures d'avant la correction, ajoutez --tune sdpa.gqa=0 --tune sdpa.split=8 --tune sdpa.pipe=0.

Les scripts de la campagne, tels qu'ils ont tourné, mes prédictions, le protocole et les données brutes : github.com/xigh/herbert-cuda-bench, tag v1-2026-10-25c. Si un chiffre vous semble faux, vous avez de quoi le vérifier.

La suite

Le prochain article passe à Metal, sur les puces d'Apple. La même question, une fois de plus : le même principe tient-il ?