Améliorations apportées à WebAssembly et WebGPU pour une IA Web plus rapide, partie 2

Ce document fait suite à la partie 1 des améliorations apportées à WebAssembly et WebGPU pour une IA Web plus rapide. Nous vous recommandons de lire cet article ou de regarder la vidéo de l'IO 24 avant de continuer.

Austin Eng
Austin Eng
Deepti Gandluri
Deepti Gandluri
François Beaufort
François Beaufort

WebGPU

WebGPU permet aux applications Web d'accéder au matériel GPU du client pour effectuer des calculs efficaces et hautement parallèles. Depuis le lancement de WebGPU dans Chrome, nous avons assisté à des démonstrations incroyables d'intelligence artificielle (IA) et de machine learning (ML) sur le Web.

Par exemple, Web Stable Diffusion a démontré qu'il était possible d'utiliser l'IA pour générer des images à partir de texte, directement dans le navigateur. Plus tôt cette année, l'équipe Mediapipe de Google a publié une fonctionnalité expérimentale pour l'inférence de grands modèles de langage.

L'animation suivante montre Gemma, le grand modèle de langage (LLM, Large Language Model) Open Source de Google, exécuté entièrement sur l'appareil dans Chrome, en temps réel.

La démonstration suivante Hugging Face du modèle Segment Anyth de Meta de Meta produit des masques d'objets de haute qualité entièrement sur le client.

Ce ne sont là que quelques exemples de projets qui mettent en avant la puissance de WebGPU pour l'IA et le ML. WebGPU permet à ces modèles et à d'autres de s'exécuter beaucoup plus rapidement qu'ils ne le pourraient sur un processeur.

Le benchmark WebGPU pour l'intégration de texte de Hugging Face montre des gains d'accélération considérables par rapport à une implémentation de processeur du même modèle. Sur un ordinateur portable Apple M1 Max, WebGPU était 30 fois plus rapide. D'autres ont rapporté que WebGPU accélère le benchmark plus de 120 fois.

Amélioration des fonctionnalités WebGPU pour l'IA et le ML

WebGPU est idéal pour les modèles d'IA et de ML, qui peuvent comporter des milliards de paramètres, grâce à la compatibilité avec les nuanceurs de calcul. Les nuanceurs de calcul s'exécutent sur le GPU et permettent d'exécuter des opérations de tableau parallèles sur de grands volumes de données.

Parmi les nombreuses améliorations apportées à WebGPU l'année dernière, nous avons continué d'ajouter des fonctionnalités pour améliorer les performances de ML et d'IA sur le Web. Nous avons récemment lancé deux nouvelles fonctionnalités: les produits à virgule flottante 16 bits et les produits empaquetés de points entiers.

Valeur à virgule flottante 16 bits

N'oubliez pas que les charges de travail de ML ne nécessitent pas de précision. shader-f16 est une fonctionnalité qui permet d'utiliser le type f16 dans le langage d'ombrage WebGPU. Ce type à virgule flottante prend 16 bits au lieu des 32 bits habituels. f16 a une plage plus petite et est moins précise, mais pour de nombreux modèles de ML, cela suffit.

Cette fonctionnalité améliore l'efficacité de plusieurs façons:

  • Mémoire réduite: les Tensors avec des éléments f16 occupent la moitié de l'espace, ce qui réduit de moitié l'utilisation de la mémoire. Les calculs GPU présentent souvent un goulot d'étranglement sur la bande passante mémoire. Par conséquent, la moitié de la mémoire peut entraîner une exécution deux fois plus rapide des nuanceurs. Techniquement, vous n'avez pas besoin de f16 pour économiser de la bande passante mémoire. Il est possible de stocker les données dans un format de faible précision, puis de les étendre à f32 complète dans le nuanceur pour effectuer des calculs. Toutefois, le GPU consomme une puissance de calcul supplémentaire pour empaqueter et décompresser les données.

  • Diminution de la conversion de données: f16 utilise moins de calculs en minimisant la conversion de données. Les données de faible précision peuvent être stockées et utilisées directement sans conversion.

  • Augmentation du parallélisme: les GPU modernes peuvent ajuster simultanément un plus grand nombre de valeurs dans leurs unités d'exécution, ce qui leur permet d'effectuer un plus grand nombre de calculs parallèles. Par exemple, un GPU acceptant jusqu'à 5 000 milliards d'opérations en virgule flottante f32 par seconde peut accepter 10 000 milliards d'opérations en virgule flottante f16 par seconde.

<ph type="x-smartling-placeholder">
</ph> Capture d&#39;écran du benchmark WebGPU pour l&#39;intégration de texte <ph type="x-smartling-placeholder">
</ph> Avec shader-f16, le benchmark WebGPU WebGPU de Hugging Face pour l'intégration de texte exécute le benchmark trois fois plus rapide que f32 sur l'ordinateur portable Apple M1 Max.

WebLLM est un projet qui peut exécuter plusieurs grands modèles de langage. Il utilise Apache TVM, un framework de compilation de machine learning Open Source.

J'ai demandé à WebLLM de planifier un voyage à Paris à l'aide du modèle Llama 3 comportant huit milliards de paramètres. Les résultats montrent que pendant la phase de préremplissage du modèle, f16 est 2,1 fois plus rapide que f32. Pendant la phase de décodage, il est plus de 1,3 fois plus rapide.

Les applications doivent d'abord confirmer que l'adaptateur GPU est compatible avec f16 et, s'il est disponible, l'activer explicitement lorsqu'il demande un appareil GPU. Si f16 n'est pas compatible, vous ne pouvez pas la demander dans le tableau requiredFeatures.

// main.js

const adapter = await navigator.gpu.requestAdapter();
const supportsF16 = adapter.features.has('shader-f16');
if (supportsF16) {
  // Use f16.
  const device = await adapter.requestDevice({
    requiredFeatures: ['shader-f16'],
  });
  initApp(device);
}

Ensuite, dans vos nuanceurs WebGPU, vous devez activer explicitement f16 en haut de la page. Vous êtes ensuite libre de l'utiliser dans le nuanceur comme n'importe quel autre type de données à virgule flottante.

// my-shader.wgsl

enable f16;

struct Data {
  values : array<vec4<f16>>
}
@group(0) @binding(0) var<storage, read> data : Data;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) {
  let value : vec4<f16> = data.values[gid.x];
  ...
}

Produits scalaires entiers conditionnés

De nombreux modèles fonctionnent toujours bien avec seulement 8 bits de précision (moitié de f16). Cette méthode est populaire parmi les LLM et les modèles d'image pour la segmentation et la reconnaissance d'objets. Cela dit, la qualité de sortie des modèles se dégrade avec une précision moindre. La quantification 8 bits ne convient donc pas à toutes les applications.

Peu de GPU sont compatibles de manière native avec les valeurs 8 bits. C'est là que les produits entier scalaires empaquetés entrent en jeu. Nous avons lancé DP4a dans Chrome 123.

Les GPU modernes ont des instructions spéciales pour prendre deux entiers de 32 bits, les interpréter comme quatre entiers de 8 bits empaquetés de manière consécutive et calculer le produit scalaire entre leurs composants.

Cela est particulièrement utile pour l'IA et le machine learning, car les noyaux de multiplication matriciel sont composés de nombreux produits scalaires.

Par exemple, multiplions une matrice 4 x 8 par un vecteur 8 x 1. Ce calcul implique de prendre des produits scalaires à quatre pour calculer chacune des valeurs du vecteur de sortie. A, B, C et D.

Schéma d&#39;un exemple de multiplication matricielle à vecteur

Le processus de calcul de ces sorties est le même : nous allons examiner les étapes du calcul de l'un d'entre eux. Avant tout calcul, nous devons d'abord convertir les données d'entiers de 8 bits dans un type avec lequel nous pouvons effectuer des calculs arithmétiques, comme f16. Ensuite, nous exécutons une multiplication au niveau des éléments et ajoutons tous les produits ensemble. Au total, pour l'ensemble de la multiplication matricielle-vecteur, nous effectuons 40 conversions en nombre entier pour flotter afin de désimbriquer les données, 32 multiplications à virgule flottante et 28 additions à virgule flottante.

Pour les matrices plus grandes avec plus d'opérations, les produits scalaires entiers empaquetés peuvent aider à réduire la quantité de travail.

Pour chacune des sorties dans le vecteur de résultats, nous effectuons deux opérations de produit scalaire regroupées à l'aide du langage dot4U8Packed intégré WebGPU Shading Language, puis nous additionnons les résultats. Au total, pour l'ensemble de la multiplication matricielle-vecteur, nous n'effectuons aucune conversion de données. Nous exécutons huit produits scalaires empaquetés et quatre additions d'entiers.

Schéma de l&#39;exemple de multiplication matricielle/vecteur entier empaqueté

Nous avons testé des produits scalaires empaquetés avec des données 8 bits sur différents GPU grand public. Par rapport à la virgule flottante 16 bits, nous pouvons voir que 8 bits est 1,6 à 2,8 fois plus rapide. Lorsque nous utilisons également des produits scalaires en nombres entiers empaquetés, les performances sont encore meilleures. Il est 1,7 à 2,9 fois plus rapide.

<ph type="x-smartling-placeholder">
</ph> Capture d&#39;écran de la multiplication matricielle à vecteur: f16 vs u8 <ph type="x-smartling-placeholder">
</ph> Graphique 1: Vitesse du vecteur matriciel, comparaison de f16 à U8 et U8 avec dot4U8Packed.

Vérifiez la compatibilité des navigateurs avec la propriété wgslLanguageFeatures. Si le GPU n'est pas compatible de manière native avec les produits scalaires empaquetés, le navigateur émule sa propre implémentation.

// main.js

if (navigator.gpu.wgslLanguageFeatures.has('packed_4x8_integer_dot_product')) {
  // Use dot4U8Packed, dot4I8Packed builtin
  // functions in the shaders.
}

La différence d'extrait de code suivante met en évidence les modifications nécessaires pour utiliser des produits entiers empaquetés dans un nuanceur WebGPU.

Avant : nuanceur WebGPU qui accumule des produits scalaires partiels dans la variable "sum". À la fin de la boucle, `sum` contient le produit scalaire complet entre un vecteur et une ligne de la matrice d'entrée.

// my-dot-product.wgsl

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
  var sum : f16;
  let start = gid.x * uniforms.dim;
  for (var i = 0u; i < uniforms.dim; i++) {
    let v1 : vec4<f16> = vector.values[i];
    let v2 : vec4<f16> = matrix.values[start + i];
    sum += dot(v1, v2);
  }
}

Après : un nuanceur WebGPU écrit pour utiliser des produits scalaires entiers empaquetés. La principale différence est qu'au lieu de charger quatre valeurs flottantes à partir du vecteur et de la matrice, ce nuanceur charge un seul entier de 32 bits. Cet entier de 32 bits contient les données de quatre valeurs entières de 8 bits. Ensuite, nous appelons dot4U8Packed pour calculer le produit scalaire des deux valeurs.

// my-dot-product.wgsl

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
  var sum : f32;
  let start = gid.x * uniforms.dim;
  for (var i = 0u; i < uniforms.dim; i++) {
    let v1 : u32 = vector.values[i];
    let v2 : u32 = matrix.values[start + i];
    sum += dot4U8Packed(v1, v2);
  }
}

Les fonctionnalités Chrome à virgule flottante 16 bits et entières empaquetées permettent d'accélérer l'IA et le ML. La virgule flottante 16 bits est disponible lorsque le matériel le permet, et Chrome implémente des produits empaquetés de type entier scalaire sur tous les appareils.

Vous pouvez dès aujourd'hui utiliser ces fonctionnalités dans la version stable de Chrome pour améliorer les performances.

Fonctionnalités proposées

À l'avenir, nous étudions deux autres caractéristiques: les sous-groupes et la multiplication matricielle coopérative.

La fonctionnalité de sous-groupes permet au parallélisme de niveau SIMD de communiquer ou d'effectuer des opérations mathématiques collectives, comme une somme pour plus de 16 nombres. Cela permet un partage de données interthread efficace. Les sous-groupes sont compatibles avec les API GPU modernes, avec des noms différents et sous des formes légèrement différentes.

Nous avons résumé l'ensemble commun dans une proposition que nous avons transmise au groupe de normalisation WebGPU. Nous avons également prototypé des sous-groupes dans Chrome à l'origine d'une fonctionnalité expérimentale, et avons inclus nos premiers résultats dans la discussion. Le principal problème est de savoir comment garantir un comportement portable.

La multiplication matricielle coopérative est un ajout plus récent aux GPU. Une grande multiplication matricielle peut être décomposée en plusieurs multiplications matricielles plus petites. La multiplication matricielle coopérative effectue des multiplications sur ces petits blocs de taille fixe en une seule étape logique. Au cours de cette étape, un groupe de threads coopère efficacement pour calculer le résultat.

Nous avons consulté la compatibilité des API GPU sous-jacentes et prévoyons de présenter une proposition au groupe de normalisation WebGPU. Comme pour les sous-groupes, nous pensons qu'une grande partie de la discussion sera axée sur la portabilité.

Pour évaluer les performances des opérations de sous-groupes, dans une application réelle, nous avons intégré la prise en charge expérimentale des sous-groupes dans MediaPipe et l'avons testée avec le prototype de Chrome pour les opérations de sous-groupes.

Nous avons utilisé des sous-groupes dans les noyaux GPU de la phase de préremplissage du grand modèle de langage. Je ne signale donc que l'accélération de la phase de préremplissage. Sur un GPU Intel, nous constatons que les sous-groupes sont deux fois et demi plus rapides que la référence. Cependant, ces améliorations ne sont pas cohérentes entre les différents GPU.

<ph type="x-smartling-placeholder">
</ph> Capture d&#39;écran de l&#39;accélération des sous-groupes dans l&#39;inférence LLM MediaPipe <ph type="x-smartling-placeholder">
</ph> Graphique 2. Grâce aux sous-groupes, le préremplissage s'exécute 2,5 fois plus vite sur le GPU Intel Tiger Lake GT2, avec une compatibilité expérimentale dans Chrome et Mediapipe.

Le graphique suivant montre les résultats de l'application de sous-groupes pour optimiser un microbenchmark de multiplication matriciel sur plusieurs GPU grand public. La multiplication matricielle est l'une des opérations les plus lourdes dans les grands modèles de langage. Les données montrent que sur de nombreux GPU, les sous-groupes augmentent la vitesse de deux, cinq, voire treize fois la vitesse de référence. Toutefois, notez que sur le premier GPU, les sous-groupes ne sont pas beaucoup mieux du tout.

<ph type="x-smartling-placeholder">
</ph> Capture d&#39;écran de l&#39;accélération du sous-groupe pour la multiplication matricielle <ph type="x-smartling-placeholder">
</ph> Graphique 3. L'application de sous-groupes pour la multiplication matricielle peut améliorer davantage les performances.

L'optimisation GPU est difficile

En fin de compte, le meilleur moyen d'optimiser votre GPU dépend du GPU proposé par le client. L'utilisation de nouvelles fonctionnalités GPU sophistiquées ne donne pas toujours les résultats escomptés, car de nombreux facteurs complexes peuvent être impliqués. La meilleure stratégie d'optimisation sur un GPU peut ne pas être la meilleure sur un autre.

Vous souhaitez réduire la bande passante mémoire tout en utilisant pleinement les threads de calcul du GPU.

Les modèles d'accès à la mémoire peuvent également être très importants. Les GPU ont tendance à offrir de bien meilleures performances lorsque les threads de calcul accèdent à la mémoire selon un modèle optimal pour le matériel. Important: Vous devez vous attendre à des caractéristiques de performances différentes selon le matériel GPU. Vous devrez peut-être exécuter différentes optimisations en fonction du GPU.

Dans le graphique suivant, nous avons pris le même algorithme de multiplication matricielle, mais ajouté une autre dimension pour démontrer davantage l'impact de différentes stratégies d'optimisation, ainsi que la complexité et la variance entre les différents GPU. Nous avons ici introduit une nouvelle technique, que nous appellerons le « Swizzle ». Swizzle optimise les modèles d'accès à la mémoire afin de les optimiser pour le matériel.

Comme vous pouvez le constater, le tourbillon de mémoire a un impact significatif, elle a parfois encore plus d'impact que les sous-groupes. Sur le GPU 6, la vitesse de swizzle est multipliée par 12, tandis que les sous-groupes offrent une vitesse 13 fois supérieure. Ensemble, ils offrent une vitesse incroyablement supérieure de 26 fois. Pour les autres GPU, il arrive que le swizzle et les sous-groupes combinés offrent de meilleures performances que l'un ou l'autre. Sur les autres GPU, c'est la technologie swizzle qui offre les meilleurs résultats.

<ph type="x-smartling-placeholder">
</ph> Capture d&#39;écran de l&#39;accélération des stratégies de multiplication matricielle <ph type="x-smartling-placeholder">
</ph> Graphique 4.

Le réglage et l'optimisation des algorithmes GPU pour qu'ils fonctionnent bien sur tous les éléments matériels peuvent nécessiter une grande expertise. Heureusement, les frameworks de bibliothèques de niveau supérieur, tels que Mediapipe, Transformers.js, Apache TVM, ONNX Runtime Web et bien d'autres font l'objet d'une énorme quantité de travail talentueux.

Les bibliothèques et les frameworks sont bien placés pour gérer la complexité liée à la gestion de diverses architectures GPU et générer du code spécifique à la plate-forme qui fonctionnera bien sur le client.

Points à retenir

L'équipe Chrome continue de contribuer à faire évoluer les normes WebAssembly et WebGPU afin d'améliorer la plate-forme Web pour les charges de travail de machine learning. Nous investissons dans des primitives de calcul plus rapides, une meilleure interopérabilité entre les normes Web, et nous nous assurons que les modèles, petits et grands, fonctionnent efficacement sur tous les appareils.

Notre objectif est de maximiser les fonctionnalités de la plate-forme tout en conservant le meilleur du Web: portée, facilité d'utilisation et portabilité. Et nous ne le faisons pas seuls. Nous travaillons en collaboration avec les autres fournisseurs de navigateurs de W3C, ainsi qu'avec de nombreux partenaires de développement.

Nous espérons que vous vous souviendrez des points suivants lorsque vous travaillez avec WebAssembly et WebGPU:

  • L'inférence IA est d'ores et déjà disponible sur le Web et sur tous les appareils. Cela présente l'avantage d'être exécuté sur des appareils clients, avec par exemple une réduction des coûts liés aux serveurs, une faible latence et une confidentialité accrue.
  • Si de nombreuses fonctionnalités présentées concernent principalement les auteurs du framework, vos applications peuvent en bénéficier sans trop de frais généraux.
  • Les normes Web sont fluides et évoluent, et nous avons toujours besoin de vos commentaires. Partagez les vôtres pour WebAssembly et WebGPU.

Remerciements

Nous souhaitons remercier l'équipe de graphisme Web d'Intel, qui a joué un rôle déterminant dans le développement de WebGPU f16 et de toutes les fonctionnalités du produit Integer Point. Nous souhaitons remercier les autres membres des groupes de travail WebAssembly et WebGPU du W3C, y compris les autres fournisseurs de navigateurs.

Merci aux équipes d'IA et de ML de Google et de la communauté Open Source pour leurs excellentes partenaires. Et bien sûr, tous nos collègues qui rendent tout cela possible.