Tutorial WebGPU: pengiraan, bucu dan pelorek serpihan di web

WebGPU ialah teknologi global yang menjanjikan untuk membawa keupayaan pengkomputeran GPU yang canggih ke web, yang memanfaatkan semua platform pengguna menggunakan pangkalan kod yang dikongsi.

Walaupun pendahulunya, WebGL, berkuasa, ia benar-benar tidak mempunyai keupayaan pengiraan shader, mengehadkan skop aplikasinya.

WGSL (WebGPU Shader/Compute Language) menggunakan amalan terbaik dari kawasan seperti Rust dan GLSL.

Semasa saya belajar menggunakan WebGPU, saya menemui beberapa jurang dalam dokumentasi: Saya berharap untuk mencari titik permulaan yang mudah untuk menggunakan pelorek pengiraan untuk mengira data bagi pelorek puncak dan serpihan.

HTML fail tunggal untuk semua kod dalam tutorial ini boleh didapati di https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb - baca terus untuk butiran terperinci.

Berikut ialah demonstrasi satu klik HTML ini yang dijalankan pada domain saya: https://www.php.cn/link/bed827b4857bf056d05980661990ccdc Penyemak imbas berasaskan WebGPU seperti Chrome atau Edge https://www.php.cn/link/bae00fb8b4115786ba5dbbb67b9b177a).

Tetapan Terperinci

Ini ialah simulasi zarah - ia berlaku mengikut langkah masa dari semasa ke semasa.

Masa dijejaki pada JS/CPU dan diserahkan kepada GPU sebagai seragam (terapung).

Data zarah diurus sepenuhnya pada GPU - walaupun masih berinteraksi dengan CPU, membenarkan memori diperuntukkan dan nilai awal ditetapkan. Ia juga mungkin untuk membaca data kembali ke CPU, tetapi ini ditinggalkan dalam tutorial ini.

Keajaiban persediaan ini ialah setiap zarah dikemas kini selari dengan semua zarah lain, membolehkan pengiraan dan kelajuan pemaparan yang luar biasa dalam penyemak imbas (penyejajaran memaksimumkan bilangan teras pada GPU; Kita boleh membahagikan bilangan zarah dengan bilangan teras untuk mendapatkan bilangan sebenar kitaran setiap langkah kemas kini setiap teras).

Ikat

Mekanisme yang digunakan WebGPU untuk pertukaran data antara CPU dan GPU adalah mengikat - Tatasusunan JS (seperti Float32Array) boleh "terikat" ke lokasi memori dalam WGSL menggunakan penimbal WebGPU. Lokasi memori WGSL dikenal pasti oleh dua integer: nombor kumpulan dan nombor pengikat.

Dalam kes kami, kedua-dua pelorek pengiraan dan pelorek puncak bergantung pada dua pengikatan data: masa dan kedudukan zarah.

Masa - pakaian seragam

Takrifan seragam wujud dalam pelorek pengiraan (https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L43) dan pelorek puncak (https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L69) Sederhana - Kira kedudukan kemas kini pelorek, warna kemas kini pelorek bucu berdasarkan masa.

Mari kita lihat persediaan mengikat dalam JS dan WGSL, bermula dengan pelorek pengiraan.

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

dan pengisytiharan yang sepadan dalam shader pengiraan

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

Yang penting, kami mengikat TimeBuffer di sebelah JS ke WGSL dengan memadankan nombor kumpulan dan nombor mengikat dalam JS dan WGSL.

Ini membolehkan kami mengawal nilai pembolehubah daripada 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>

Kedudukan Zarah - Storan WGSL

Kami menyimpan dan mengemas kini kedudukan zarah terus dalam memori yang boleh diakses GPU – membolehkan kami mengemas kininya secara selari dengan memanfaatkan seni bina berbilang teras besar GPU.

Persejajaran diselaraskan dengan bantuan saiz kumpulan kerja, diisytiharkan dalam lorek pengiraan:

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

@builtin(global_invocation_id) global_id : vec3 Nilai menyediakan pengecam urutan.

Mengikut takrifan, global_invocation_id = workgroup_id * workgroup_size local_invocation_id - ini bermakna ia boleh digunakan sebagai indeks zarah.

Sebagai contoh, jika kita mempunyai 10k zarah dan saiz kumpulan kerja ialah 64, kita perlu menjadualkan kumpulan kerja Math.ceil(10000/64). Setiap kali pas pengiraan dicetuskan daripada JS, kami akan secara jelas memberitahu GPU untuk melaksanakan jumlah kerja ini:

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

Jika PARTICLE_COUNT == 10000 dan WORKGROUP_SIZE == 64, kami akan memulakan 157 kumpulan kerja (10000/64 = 156.25), dan julat pengiraan local_invocation_id bagi setiap kumpulan kerja ialah 0 hingga 63 (manakala julat kumpulan_kerja_15 ialah ). Oleh kerana 157 * 64 = 1048, kami akhirnya akan melakukan lebih sedikit pengiraan dalam kumpulan kerja. Kami mengendalikan limpahan dengan membuang panggilan berlebihan.

Berikut ialah keputusan akhir pengiraan shader selepas mengambil kira faktor ini:

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

Nilai ini akan kekal merentasi pas pengiraan kerana zarah ditakrifkan sebagai pembolehubah storan.

Baca kedudukan zarah dalam lorek pengiraan dalam lorek bucu

Untuk membaca kedudukan zarah dalam pelorek puncak daripada pelorek pengiraan, kita memerlukan paparan baca sahaja, kerana hanya pelorek pengiraan boleh menulis ke storan.

Berikut ialah kenyataan daripada 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>

Cuba untuk menggunakan semula gaya baca_tulis yang sama dalam pelorek pengiraan hanya akan ralat:

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

Ambil perhatian bahawa nombor pengikatan dalam pelorek puncak tidak perlu sepadan dengan pengiraan nombor pengikatan pelorek - nombor itu hanya perlu sepadan dengan pengisytiharan kumpulan pengikat pelorek puncak:

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

Saya memilih binding:2 dalam kod sampel GitHub https://www.php.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L70 - hanya untuk meneroka sempadan kekangan Web yang dikenakan oleh GPU

Jalankan simulasi langkah demi langkah

Dengan semua tetapan disediakan, kemas kini dan gelung pemaparan diselaraskan dalam 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>

Kesimpulan

WebGPU mengeluarkan kuasa pengkomputeran GPU selari secara besar-besaran dalam penyemak imbas.

Ia berjalan dalam pas - setiap pas mempunyai pembolehubah setempat yang didayakan melalui saluran paip dengan pengikatan memori (merapatkan memori CPU dan memori GPU).

Penghantaran pengiraan membolehkan penyelarasan beban kerja selari melalui kumpulan kerja.

Walaupun ia memerlukan beberapa persediaan yang berat, saya rasa gaya pengikatan/keadaan tempatan adalah peningkatan yang besar berbanding model keadaan global WebGL - menjadikannya lebih mudah untuk digunakan sambil akhirnya membawa kuasa pengkomputeran GPU ke Entered the Web.

Atas ialah kandungan terperinci Tutorial WebGPU: pengiraan, bucu dan pelorek serpihan di web. Untuk maklumat lanjut, sila ikut artikel berkaitan lain di laman web China PHP!