Cet article explore l'API WebGPU expérimentale à travers des exemples et aide vous commencez à effectuer des calculs avec parallélisme des données à l'aide du GPU.
Contexte
Comme vous le savez peut-être déjà, le processeur graphique (GPU) est un processeur électronique sous-système d’un ordinateur qui était à l’origine spécialisé dans le traitement graphiques. Cependant, au cours des 10 dernières années, elle a évolué vers un modèle plus flexible permettant aux développeurs d'implémenter de nombreux types d'algorithmes, pas seulement des graphismes en 3D tout en profitant de l'architecture unique GPU. C'est ce que l'on appelle le calcul GPU. Un GPU est utilisé un coprocesseur d'informatique scientifique d'usage général est appelé "à usage général" de la programmation GPU (GPGPU).
Le calcul GPU a grandement contribué au boom récent du machine learning, car les réseaux de neurones à convolution et d'autres modèles peuvent exploiter pour s'exécuter plus efficacement sur les GPU. Avec la plate-forme Web actuelle manque de capacités de calcul GPU, le "GPU pour le Web" du W3C Groupe de la communauté conçoit une API pour exposer les API GPU modernes disponibles sur la plupart des appareils actuels. Cette API s'appelle WebGPU.
WebGPU est une API de bas niveau, comme WebGL. Il est très puissant et assez détaillé, comme vous le verrez. Mais ce n'est pas grave. Ce que nous recherchons, ce sont des performances.
Dans cet article, nous allons nous concentrer sur la partie "GPU Compute" de WebGPU, et pour être je ne fais qu'effleurer la surface, pour que vous puissiez commencer à jouer vous-même. Je vais m'intéresser de plus près au rendu WebGPU (canevas, texture, etc.) dans les articles à venir.
Accéder au GPU
L'accès au GPU est facile dans WebGPU. Vous appelez navigator.gpu.requestAdapter()...
renvoie une promesse JavaScript qui sera résolue de manière asynchrone avec un GPU
adaptateur secteur. Considérez cet adaptateur comme une carte graphique. Il peut être intégré
(sur la même puce que le CPU) ou discrètes (généralement une carte PCIe plus
performante mais qui utilise plus de puissance).
Une fois que vous disposez de l'adaptateur GPU, appelez adapter.requestDevice() pour obtenir une promesse
qui se résoudre avec un périphérique GPU que vous utiliserez
pour effectuer des calculs GPU.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
Les deux fonctions proposent des options qui vous permettent d'être précis sur le type l'adaptateur secteur (préférence d'alimentation) et l'appareil (extensions, limites) souhaités. Pour le Par souci de simplicité, nous utiliserons les options par défaut de cet article.
Écriture de la mémoire tampon
Voyons comment utiliser JavaScript pour écrire des données en mémoire pour le GPU. Ce n'est pas simple à cause du modèle de bac à sable utilisé dans les environnements des navigateurs.
L'exemple ci-dessous montre comment écrire quatre octets dans la mémoire tampon accessible
à partir du GPU. Elle appelle device.createBuffer(), qui prend la taille du
tampon et son utilisation. Même si l'indicateur d'utilisation GPUBufferUsage.MAP_WRITE est
ne sont pas nécessaires pour cet appel. Précisons que nous voulons écrire
dans ce tampon. Il en résulte un objet de tampon GPU mappé lors de la création grâce à
mappedAtCreation défini sur "true". Le tampon de données binaires brutes associé peut alors
être récupérées en appelant la méthode de tampon GPU getMappedRange().
L'écriture d'octets est un processus courant si vous avez déjà joué avec ArrayBuffer. utilisez un
TypedArray et copiez-y les valeurs.
// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuBuffer = device.createBuffer({
mappedAtCreation: true,
size: 4,
usage: GPUBufferUsage.MAP_WRITE
});
const arrayBuffer = gpuBuffer.getMappedRange();
// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);
À ce stade, le tampon GPU est mappé, ce qui signifie qu'il appartient au processeur.
il est accessible en lecture/écriture
à partir de JavaScript. Pour que le GPU puisse y accéder,
doit être supprimé, ce qui est aussi simple que d'appeler gpuBuffer.unmap().
Le concept de mappage/non mappé est nécessaire pour éviter les conditions de concurrence où les GPU et la mémoire d'accès du processeur en même temps.
Lecture de la mémoire tampon
Voyons maintenant comment copier un tampon GPU dans un autre tampon GPU et comment le relire.
Comme nous écrivons dans le premier tampon GPU et que nous voulons le copier dans un second
Tampon GPU, un nouvel indicateur d'utilisation GPUBufferUsage.COPY_SRC est requis. Le deuxième
Cette fois, le tampon GPU est créé dans un état non mappé avec
device.createBuffer() Son indicateur d'utilisation est GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ, car il sera utilisé comme destination du premier GPU
tampon et lu en JavaScript une fois que les commandes de copie GPU ont été exécutées.
// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuWriteBuffer = device.createBuffer({
mappedAtCreation: true,
size: 4,
usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.COPY_SRC
});
const arrayBuffer = gpuWriteBuffer.getMappedRange();
// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);
// Unmap buffer so that it can be used later for copy.
gpuWriteBuffer.unmap();
// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
size: 4,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});
Comme le GPU est un coprocesseur indépendant, toutes les commandes GPU sont exécutées.
de manière asynchrone. C'est pourquoi il existe une liste de commandes
GPU créées et envoyées
par lot selon les besoins. Dans WebGPU, l'encodeur de commande GPU renvoyé par
device.createCommandEncoder() est l'objet JavaScript qui crée un lot
"en mémoire tampon" qui seront envoyées au GPU à un moment donné. Les méthodes sur
Les GPUBuffer, en revanche, sont "sans mémoire tampon", ce qui signifie qu'elles s'exécutent de manière atomique
au moment où elles sont appelées.
Une fois que vous disposez de l'encodeur de commande GPU, appelez copyEncoder.copyBufferToBuffer().
comme indiqué ci-dessous pour ajouter cette commande à la file d'attente de commandes pour une exécution ultérieure.
Enfin, terminez les commandes d'encodage en appelant copyEncoder.finish() et en envoyant
à la file d'attente de commandes
de l'appareil GPU. La file d'attente est chargée de gérer
les envois effectués via device.queue.submit() avec les commandes GPU en tant qu'arguments.
Cela exécutera de manière atomique toutes les commandes stockées dans le tableau dans l'ordre.
// Encode commands for copying buffer to buffer.
const copyEncoder = device.createCommandEncoder();
copyEncoder.copyBufferToBuffer(
gpuWriteBuffer /* source buffer */,
0 /* source offset */,
gpuReadBuffer /* destination buffer */,
0 /* destination offset */,
4 /* size */
);
// Submit copy commands.
const copyCommands = copyEncoder.finish();
device.queue.submit([copyCommands]);
À ce stade, les commandes de file d'attente GPU ont été envoyées, mais pas nécessairement exécutées.
Pour lire le deuxième tampon GPU, appelez gpuReadBuffer.mapAsync() avec
GPUMapMode.READ Elle renvoie une promesse qui se résoudre lorsque le tampon GPU sera
mappés. Obtenez ensuite la plage mappée avec gpuReadBuffer.getMappedRange(),
contient les mêmes valeurs que le premier tampon GPU une fois que toutes les commandes GPU mises en file d'attente
ont été exécutées.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));
Vous pouvez essayer cet exemple.
En bref, voici ce que vous devez retenir concernant les opérations de la mémoire tampon:
- Vous devez annuler le mappage des tampons GPU pour pouvoir les utiliser dans la file d'attente de l'appareil.
- Lorsqu'ils sont mappés, les tampons GPU peuvent être lus et écrits en JavaScript.
- Les tampons GPU sont mappés lorsque
mapAsync()etcreateBuffer()avecmappedAtCreationdéfinies sur "true" sont appelées.
Programmation du nuanceur
Programmes s'exécutant sur le GPU qui n'effectuent que des calculs (et ne dessinent pas des triangles) sont appelés nuanceurs de calcul. Elles sont exécutées en parallèle par des centaines de cœurs de GPU (plus petits que les cœurs de CPU) qui fonctionnent ensemble pour traiter données. Leurs entrées et sorties sont des tampons dans WebGPU.
Pour illustrer l'utilisation des nuanceurs de calcul dans WebGPU, nous allons jouer avec la matrice la multiplication, un algorithme de machine learning couramment utilisé et illustré ci-dessous.
<ph type="x-smartling-placeholder">
En résumé, voici ce que nous allons faire:
- Créez trois tampons GPU (deux pour les matrices à multiplier et un pour la matrice de résultats)
- Décrire les entrées et les sorties du nuanceur de calcul
- Compiler le code du nuanceur de calcul
- Configurer un pipeline de calcul
- Envoyer par lots les commandes encodées au GPU
- Lire le tampon GPU de la matrice des résultats
Création de tampons GPU
Par souci de simplicité, les matrices seront représentées sous la forme d'une liste de valeurs flottantes les nombres de points. Le premier élément est le nombre de lignes, tandis que le deuxième élément est le nombre de colonnes, et le reste correspond aux nombres réels de la matrice.
<ph type="x-smartling-placeholder">
Les trois tampons GPU sont des tampons de stockage, car nous devons stocker et récupérer des données dans
le nuanceur de calcul. Cela explique pourquoi les indicateurs d'utilisation de la mémoire tampon GPU incluent
GPUBufferUsage.STORAGE pour l'ensemble d'entre elles. L'indicateur d'utilisation de la matrice de résultats comporte également
GPUBufferUsage.COPY_SRC, car elle sera copiée dans un autre tampon pour
une fois que toutes les commandes
de file d'attente GPU ont été exécutées.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
// First Matrix
const firstMatrix = new Float32Array([
2 /* rows */, 4 /* columns */,
1, 2, 3, 4,
5, 6, 7, 8
]);
const gpuBufferFirstMatrix = device.createBuffer({
mappedAtCreation: true,
size: firstMatrix.byteLength,
usage: GPUBufferUsage.STORAGE,
});
const arrayBufferFirstMatrix = gpuBufferFirstMatrix.getMappedRange();
new Float32Array(arrayBufferFirstMatrix).set(firstMatrix);
gpuBufferFirstMatrix.unmap();
// Second Matrix
const secondMatrix = new Float32Array([
4 /* rows */, 2 /* columns */,
1, 2,
3, 4,
5, 6,
7, 8
]);
const gpuBufferSecondMatrix = device.createBuffer({
mappedAtCreation: true,
size: secondMatrix.byteLength,
usage: GPUBufferUsage.STORAGE,
});
const arrayBufferSecondMatrix = gpuBufferSecondMatrix.getMappedRange();
new Float32Array(arrayBufferSecondMatrix).set(secondMatrix);
gpuBufferSecondMatrix.unmap();
// Result Matrix
const resultMatrixBufferSize = Float32Array.BYTES_PER_ELEMENT * (2 + firstMatrix[0] * secondMatrix[1]);
const resultMatrixBuffer = device.createBuffer({
size: resultMatrixBufferSize,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
});
Lier la mise en page du groupe et le groupe de liaisons
Les concepts de mise en page et de groupe de liaisons sont spécifiques à WebGPU. Une liaison la mise en page de groupe définit l'interface d'entrée/sortie attendue par un nuanceur, tandis qu'un un groupe de liaisons représente les données d'entrée/sortie réelles d'un nuanceur.
Dans l'exemple ci-dessous, la mise en page du groupe de liaisons attend deux tampons de stockage en lecture seule
des liaisons d'entrée numérotées 0, 1 et un tampon de stockage à 2 pour le nuanceur de calcul.
Le groupe de liaisons, défini pour cette mise en page, associe
Le GPU met en mémoire tampon les entrées: gpuBufferFirstMatrix à la liaison 0,
gpuBufferSecondMatrix à la liaison 1 et resultMatrixBuffer à la liaison
2.
const bindGroupLayout = device.createBindGroupLayout({
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: "read-only-storage"
}
},
{
binding: 1,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: "read-only-storage"
}
},
{
binding: 2,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: "storage"
}
}
]
});
const bindGroup = device.createBindGroup({
layout: bindGroupLayout,
entries: [
{
binding: 0,
resource: {
buffer: gpuBufferFirstMatrix
}
},
{
binding: 1,
resource: {
buffer: gpuBufferSecondMatrix
}
},
{
binding: 2,
resource: {
buffer: resultMatrixBuffer
}
}
]
});
Calculer le code du nuanceur
Le code du nuanceur de calcul permettant de multiplier les matrices est écrit en WGSL, la
Le langage de nuanceur WebGPU, qui peut être traduit en SPIR-V de manière triviale. Sans
vous trouverez ci-dessous les trois tampons de stockage identifiés
avec var<storage>. Le programme utilisera firstMatrix et secondMatrix comme
entrées et resultMatrix comme sortie.
Notez que chaque tampon de stockage dispose d'une décoration binding qui correspond à
le même index défini dans les mises en page de groupes de liaisons et les groupes de liaisons déclarés ci-dessus.
const shaderModule = device.createShaderModule({
code: `
struct Matrix {
size : vec2f,
numbers: array<f32>,
}
@group(0) @binding(0) var<storage, read> firstMatrix : Matrix;
@group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
@group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;
@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) global_id : vec3u) {
// Guard against out-of-bounds work group sizes
if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
return;
}
resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);
let resultCell = vec2(global_id.x, global_id.y);
var result = 0.0;
for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
let a = i + resultCell.x * u32(firstMatrix.size.y);
let b = resultCell.y + i * u32(secondMatrix.size.y);
result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
}
let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
resultMatrix.numbers[index] = result;
}
`
});
Configuration du pipeline
Le pipeline de calcul est l'objet qui décrit réellement l'opération de calcul
que nous allons exécuter. Créez-la en appelant device.createComputePipeline().
Elle utilise deux arguments: la mise en page de groupe de liaisons que nous avons créée précédemment et un argument de calcul
Étape définissant le point d'entrée de notre nuanceur de calcul (la fonction WGSL main)
et le module de nuanceur de calcul créé avec device.createShaderModule().
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
Envoi de commandes
Après avoir instancié un groupe de liaisons avec nos trois tampons GPU et une couche de calcul avec une mise en page de groupe de liaisons, il est temps de les utiliser.
Commençons par un encodeur programmable d'une passe de calcul
commandEncoder.beginComputePass() Nous l'utiliserons pour encoder les commandes GPU
qui effectuera la multiplication de matrices. Définissez son pipeline avec
passEncoder.setPipeline(computePipeline) et son groupe de liaisons à l'index 0 avec
passEncoder.setBindGroup(0, bindGroup) L'indice 0 correspond aux
Décoration group(0) dans le code WGSL.
Voyons maintenant comment ce nuanceur de calcul va s'exécuter sur le GPU. Notre
l'objectif est d'exécuter ce programme en parallèle
pour chaque cellule de la matrice de résultats,
étape par étape. Pour une matrice de résultats de taille 16 x 32, par exemple, pour encoder
la commande d'exécution, sur un @workgroup_size(8, 8), nous l'appellerons
passEncoder.dispatchWorkgroups(2, 4) ou passEncoder.dispatchWorkgroups(16 / 8, 32 / 8).
Le premier argument "x" est la première dimension, la deuxième "y" est la deuxième dimension,
et le dernier "z" est la troisième dimension qui est définie par défaut sur 1, car nous n'en avons pas besoin ici.
Dans l'environnement de calcul GPU, l'encodage d'une commande permettant d'exécuter une fonction de noyau sur un ensemble de données s'appelle "distribution".
La taille de la grille du groupe de travail pour notre nuanceur de calcul est (8, 8) dans notre WGSL
du code source. Pour cette raison, « x » et "y" qui représentent respectivement le nombre de lignes
la première matrice et le nombre de colonnes de la deuxième matrice seront divisés
par 8. Nous pouvons maintenant envoyer un appel de calcul avec
passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) La
le nombre de grilles de groupes de travail à exécuter correspond aux arguments dispatchWorkgroups().
Comme le montre le schéma ci-dessus, chaque nuanceur aura accès à une
Objet builtin(global_invocation_id) qui sera utilisé pour savoir quel résultat
une cellule matricielle à calculer.
const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
const workgroupCountX = Math.ceil(firstMatrix[0] / 8);
const workgroupCountY = Math.ceil(secondMatrix[1] / 8);
passEncoder.dispatchWorkgroups(workgroupCountX, workgroupCountY);
passEncoder.end();
Pour arrêter l'encodeur de carte de calcul, appelez passEncoder.end(). Ensuite, créez un
Tampon GPU à utiliser comme destination pour copier le tampon de la matrice des résultats avec
copyBufferToBuffer Enfin, terminez les commandes d'encodage avec
copyEncoder.finish(), puis les envoyer à la file d'attente des appareils GPU en appelant
device.queue.submit() par les commandes GPU.
// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
size: resultMatrixBufferSize,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});
// Encode commands for copying buffer to buffer.
commandEncoder.copyBufferToBuffer(
resultMatrixBuffer /* source buffer */,
0 /* source offset */,
gpuReadBuffer /* destination buffer */,
0 /* destination offset */,
resultMatrixBufferSize /* size */
);
// Submit GPU commands.
const gpuCommands = commandEncoder.finish();
device.queue.submit([gpuCommands]);
Lire la matrice des résultats
Lire la matrice de résultats est aussi simple que d'appeler gpuReadBuffer.mapAsync() avec
GPUMapMode.READ et en attendant la résolution de la promesse de retour, ce qui indique
le tampon du GPU est mappé. À ce stade, il est possible de mapper
plage avec gpuReadBuffer.getMappedRange().
Dans notre code, le résultat enregistré dans la console JavaScript des outils de développement est "2, 2, 50, 60, 114, 140".
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));
Félicitations ! C'est fait ! Vous pouvez écouter le Sample.
Une dernière astuce
Pour rendre votre code plus facile à lire, vous pouvez utiliser la fonction
La méthode getBindGroupLayout du pipeline de calcul pour déduire le groupe de liaisons
à partir du module du nuanceur. Grâce à cette astuce, il n'est plus nécessaire de créer
une mise en page de groupe de liaisons personnalisée et spécifier une mise en page de pipeline dans vos ressources de calcul
comme illustré ci-dessous.
Une illustration de getBindGroupLayout pour l'exemple précédent est disponible.
const computePipeline = device.createComputePipeline({
- layout: device.createPipelineLayout({
- bindGroupLayouts: [bindGroupLayout]
- }),
compute: {
-// Bind group layout and bind group
- const bindGroupLayout = device.createBindGroupLayout({
- entries: [
- {
- binding: 0,
- visibility: GPUShaderStage.COMPUTE,
- buffer: {
- type: "read-only-storage"
- }
- },
- {
- binding: 1,
- visibility: GPUShaderStage.COMPUTE,
- buffer: {
- type: "read-only-storage"
- }
- },
- {
- binding: 2,
- visibility: GPUShaderStage.COMPUTE,
- buffer: {
- type: "storage"
- }
- }
- ]
- });
+// Bind group
const bindGroup = device.createBindGroup({
- layout: bindGroupLayout,
+ layout: computePipeline.getBindGroupLayout(0 /* index */),
entries: [
Résultats concernant les performances
Quelle est la différence entre l'exécution de la multiplication matricielle sur un GPU et l'exécution sur un Processeur ? Pour le savoir, j'ai écrit le programme qui vient d'être décrit pour un CPU. Comme vous le pouvez, comme le montre le graphique ci-dessous, utiliser toute la puissance du GPU semble être un choix évident lorsque la taille des matrices est supérieure à 256 x 256.
<ph type="x-smartling-placeholder">
Cet article n'était que le début de mon parcours d'exploration de WebGPU. Attendez-vous à plus Les articles présentent bientôt plus en détail le calcul GPU et comment le rendu (canevas, texture, échantillonneur) fonctionne dans WebGPU.