Tutoriel WebGPU : shaders de calcul, de sommets et de fragments sur le Web

WebGPU est une technologie mondiale qui promet d'apporter des capacités informatiques GPU de pointe au Web, bénéficiant à toutes les plates-formes grand public utilisant une base de code partagée.

Bien que son prédécesseur, WebGL, soit puissant, il manque sérieusement de capacités de calcul de shader, ce qui limite sa portée d'application.

WGSL (WebGPU Shader/Compute Language) s'appuie sur les meilleures pratiques de domaines comme Rust et GLSL.

Alors que j'apprenais à utiliser WebGPU, je suis tombé sur quelques lacunes dans la documentation : j'espérais trouver un point de départ simple pour utiliser les computing shaders afin de calculer les données des vertex et des fragment shaders.

Le fichier HTML unique pour tout le code de ce tutoriel peut être trouvé sur https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb - lisez la suite pour une ventilation détaillée.

Voici une démonstration en un seul clic de ce HTML exécuté sur mon domaine : https://www.php.cn/link/bed827b4857bf056d05980661990ccdc Un navigateur basé sur WebGPU tel que Chrome ou Edge https://www.php.cn/link/bae00fb8b4115786ba5dbbb67b9b177a).

Paramètres avancés

Il s'agit d'une simulation de particules - cela se produit par intervalles de temps.

Le temps est suivi sur JS/CPU et transmis au GPU comme uniforme (flottant).

Les données de particules sont entièrement gérées sur le GPU - bien qu'elles interagissent toujours avec le CPU, permettant d'allouer de la mémoire et de définir des valeurs initiales. Il est également possible de relire les données vers le CPU, mais ceci est omis dans ce didacticiel.

La magie de cette configuration est que chaque particule est mise à jour en parallèle avec toutes les autres particules, ce qui permet des vitesses de calcul et de rendu incroyables dans le navigateur (la parallélisation maximise le nombre de cœurs sur le GPU ; on peut diviser le nombre de particules par le nombre de cœurs pour obtenir le nombre réel de cycles par étape de mise à jour par cœur).

Relier

Le mécanisme utilisé par WebGPU pour l'échange de données entre le CPU et le GPU est contraignant - les tableaux JS (tels que Float32Array) peuvent être "liés" aux emplacements de mémoire dans WGSL à l'aide des tampons WebGPU. Les emplacements mémoire WGSL sont identifiés par deux entiers : le numéro de groupe et le numéro de liaison.

Dans notre cas, le computing shader et le vertex shader s'appuient sur deux liaisons de données : le temps et la position des particules.

Temps - uniformes

Des définitions uniformes existent dans les computing shaders (https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L43) et les vertex shaders (https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L69) Moyen - Calcule la position de mise à jour du shader, le vertex shader met à jour la couleur en fonction du temps.

Jetons un coup d'œil à la configuration des liaisons dans JS et WGSL, en commençant par les shaders de calcul.

<code>const computeBindGroup = device.createBindGroup({
  /*
    参见 computePipeline 定义，网址为
    https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L102

    它允许将 JS 字符串与 WGSL 代码链接到 WebGPU
  */
  layout: computePipeline.getBindGroupLayout(0), // 组号 0
  entries: [{
    // 时间绑定在绑定号 0
    binding: 0,
    resource: {
      /*
      作为参考，缓冲区声明为：

      const timeBuffer = device.createBuffer({
        size: Float32Array.BYTES_PER_ELEMENT,
        usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST})
      })

      https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L129
      */
      buffer: timeBuffer
    }
  },
  {
    // 粒子位置数据在绑定号 1（仍在组 0）
    binding: 1,
    resource: {
      buffer: particleBuffer
    }
  }]
});</code>

et la déclaration correspondante dans le computing shader

<code>// 来自计算着色器 - 顶点着色器中也有类似的声明
@group(0) @binding(0) var<uniform> t: f32;
@group(0) @binding(1) var<storage read_write=""> particles : array<particle>;
</particle></storage></uniform></code>

Il est important de noter que nous lions le timeBuffer du côté JS à WGSL en faisant correspondre le numéro de groupe et le numéro de liaison dans JS et WGSL.

Cela nous permet de contrôler la valeur de la variable depuis JS :

<code>/* 数组中只需要 1 个元素，因为时间是单个浮点值 */
const timeJs = new Float32Array(1)
let t = 5.3
/* 纯 JS，只需设置值 */
timeJs.set([t], 0)
/* 将数据从 CPU/JS 传递到 GPU/WGSL */
device.queue.writeBuffer(timeBuffer, 0, timeJs);</code>

Position des particules - Stockage WGSL

Nous stockons et mettons à jour les positions des particules directement dans la mémoire accessible par le GPU, ce qui nous permet de les mettre à jour en parallèle en tirant parti de l'architecture multicœur massive du GPU.

La parallélisation est coordonnée à l'aide de la taille du groupe de travail, déclarée dans le computing shader :

<code>@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
  // ...
}
</u32></code>

@builtin(global_invocation_id) global_id : vec3 La valeur fournit l'identifiant du thread.

Par définition, global_invocation_id = workgroup_id * workgroup_size local_invocation_id - cela signifie qu'il peut être utilisé comme index de particules.

Par exemple, si nous avons 10 000 particules et que workgroup_size est de 64, nous devons planifier des groupes de travail Math.ceil(10000/64). Chaque fois qu'une passe de calcul est déclenchée depuis JS, nous indiquerons explicitement au GPU d'effectuer cette quantité de travail :

<code>computePass.dispatchWorkgroups(Math.ceil(PARTICLE_COUNT / WORKGROUP_SIZE));</code>

Si PARTICLE_COUNT == 10000 et WORKGROUP_SIZE == 64, nous démarrerons 157 groupes de travail (10000/64 = 156,25), et la plage calculée de local_invocation_id de chaque groupe de travail est de 0 à 63 (tandis que la plage de workgroup_id est de 0 à 157 ). Puisque 157 * 64 = 1048, nous finirons par faire un peu plus de calculs en groupe de travail. Nous gérons les débordements en éliminant les appels redondants.

Voici le résultat final du calcul du shader après prise en compte de ces facteurs :

<code>@compute @workgroup_size(${WORKGROUP_SIZE})
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
  let index = global_id.x;
  // 由于工作组网格未对齐，因此丢弃额外的计算
  if (index >= arrayLength(&particles)) {
    return;
  }
  /* 将整数索引转换为浮点数，以便我们可以根据索引（和时间）计算位置更新 */
  let fi = f32(index);
  particles[index].position = vec2<f32>(
    /* 公式背后没有宏伟的意图 - 只不过是用时间+索引的例子 */
    cos(fi * 0.11) * 0.8 + sin((t + fi)/100)/10,
    sin(fi * 0.11) * 0.8 + cos((t + fi)/100)/10
  );
}
</f32></u32></code>

Ces valeurs persisteront à travers les passes de calcul car les particules sont définies comme variables de stockage.

Lire la position des particules dans le calculate shader dans le vertex shader

Afin de lire les positions des particules dans le vertex shader à partir du Compute Shader, nous avons besoin d'une vue en lecture seule, puisque seul le Compute Shader peut écrire dans le stockage.

Ce qui suit est une déclaration de WGSL :

<code>@group(0) @binding(0) var<uniform> t: f32;
@group(0) @binding(1) var<storage> particles : array<vec2>>;
/*
或等效：

@group(0) @binding(1) var<storage read=""> particles : array<vec2>>;
*/
</vec2></storage></vec2></storage></uniform></code>

Essayer de réutiliser le même style read_write dans un shader de calcul entraînera simplement une erreur :

<code>var with 'storage' address space and 'read_write' access mode cannot be used by vertex pipeline stage</code>

Notez que les numéros de liaison dans le vertex shader ne doivent pas nécessairement correspondre aux numéros de liaison du calculate shader - ils doivent seulement correspondre à la déclaration du groupe de liaison du vertex shader :

<code>const renderBindGroup = device.createBindGroup({
  layout: pipeline.getBindGroupLayout(0),
  entries: [{
    binding: 0,
    resource: {
      buffer: timeBuffer
    }
  },
  {
    binding: 1,
    resource: {
      buffer: particleBuffer
    }
  }]
});</code>

J'ai sélectionné contraignant:2 dans l'exemple de code GitHub https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L70 - juste pour explorer les limites des contraintes imposées par WebGPU

Lancer la simulation étape par étape

Une fois tous les paramètres en place, les boucles de mise à jour et de rendu sont coordonnées en JS :

<code>/* 从 t = 0 开始模拟 */
let t = 0
function frame() {
  /*
    为简单起见，使用恒定整数时间步 - 无论帧速率如何，都会一致渲染。
  */
  t += 1
  timeJs.set([t], 0)
  device.queue.writeBuffer(timeBuffer, 0, timeJs);

  // 计算传递以更新粒子位置
  const computePassEncoder = device.createCommandEncoder();
  const computePass = computePassEncoder.beginComputePass();
  computePass.setPipeline(computePipeline);
  computePass.setBindGroup(0, computeBindGroup);
  // 重要的是要调度正确数量的工作组以处理所有粒子
  computePass.dispatchWorkgroups(Math.ceil(PARTICLE_COUNT / WORKGROUP_SIZE));
  computePass.end();
  device.queue.submit([computePassEncoder.finish()]);

  // 渲染传递
  const commandEncoder = device.createCommandEncoder();
  const passEncoder = commandEncoder.beginRenderPass({
    colorAttachments: [{
      view: context.getCurrentTexture().createView(),
      clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 },
      loadOp: 'clear',
      storeOp: 'store',
    }]
  });
  passEncoder.setPipeline(pipeline);
  passEncoder.setBindGroup(0, renderBindGroup);
  passEncoder.draw(PARTICLE_COUNT);
  passEncoder.end();
  device.queue.submit([commandEncoder.finish()]);

  requestAnimationFrame(frame);
}
frame();</code>

Conclusion

WebGPU libère la puissance du calcul GPU massivement parallèle dans le navigateur.

Il s'exécute par passes - chaque passe a des variables locales activées via un pipeline avec liaison de mémoire (reliant la mémoire CPU et la mémoire GPU).

La livraison de calcul permet la coordination de charges de travail parallèles via des groupes de travail.

Bien que cela nécessite une configuration lourde, je pense que le style de liaison/état local est une énorme amélioration par rapport au modèle d'état global de WebGL - le rendant plus facile à utiliser tout en apportant enfin la puissance du calcul GPU pour accéder au Web.

Ce qui précède est le contenu détaillé de. pour plus d'informations, suivez d'autres articles connexes sur le site Web de PHP en chinois!