WebGPU is a global technology that promises to bring cutting-edge GPU computing capabilities to the web, benefiting all consumer platforms using a shared code base.

Although its predecessor, WebGL, is powerful, it seriously lacks compute shader capabilities, limiting its application scope.

WGSL (WebGPU Shader/Compute Language) draws on best practices from areas like Rust and GLSL.

As I was learning to use WebGPU, I came across some gaps in the documentation: I was hoping to find a simple starting point for using compute shaders to compute data for vertex and fragment shaders.

The single-file HTML for all the code in this tutorial can be found at http://m.miracleart.cn/link/2e5281ee978b78d6f5728aad8f28fedb - read on for a detailed breakdown.

Here is a single-click demonstration of this HTML running on my domain: http://m.miracleart.cn/link/bed827b4857bf056d05980661990ccdc A WebGPU-based browser such as Chrome or Edge http://m.miracleart.cn/link/bae00fb8b4115786ba5dbbb67b9b177a).

Advanced Settings

This is a particle simulation - it happens in time steps over time.

Time is tracked on JS/CPU and passed to GPU as (float) uniform.

Particle data is managed entirely on the GPU - although still interacting with the CPU, allowing memory to be allocated and initial values ??set. It is also possible to read the data back to the CPU, but this is omitted in this tutorial.

The magic of this setup is that each particle is updated in parallel with all other particles, allowing for incredible calculation and rendering speeds in the browser (parallelization maximizes the number of cores on the GPU; We can divide the number of particles by the number of cores to get the true number of cycles per update step per core).

Bind

The mechanism WebGPU uses for data exchange between CPU and GPU is binding - JS arrays (such as Float32Array) can be "bound" to memory locations in WGSL using WebGPU buffers. WGSL memory locations are identified by two integers: the group number and the binding number.

In our case, both the compute shader and the vertex shader rely on two data bindings: time and particle position.

Time - uniforms

Uniform definitions exist in compute shaders (http://m.miracleart.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L43) and vertex shaders (http://m.miracleart.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L69) Medium - Calculate shader update position, vertex shader updates color based on time.

Let’s take a look at the binding setup in JS and WGSL, starting with compute shaders.

<code>const computeBindGroup = device.createBindGroup({
  /*
    參見 computePipeline 定義，網(wǎng)址為
    http://m.miracleart.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L102

    它允許將 JS 字符串與 WGSL 代碼鏈接到 WebGPU
  */
  layout: computePipeline.getBindGroupLayout(0), // 組號 0
  entries: [{
    // 時間綁定在綁定號 0
    binding: 0,
    resource: {
      /*
      作為參考，緩沖區(qū)聲明為：

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

      http://m.miracleart.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L129
      */
      buffer: timeBuffer
    }
  },
  {
    // 粒子位置數(shù)據(jù)在綁定號 1（仍在組 0）
    binding: 1,
    resource: {
      buffer: particleBuffer
    }
  }]
});</code>

and the corresponding declaration in the compute 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>

Importantly, we bind the timeBuffer on the JS side to WGSL by matching the group number and binding number in JS and WGSL.

This allows us to control the value of the variable from JS:

<code>/* 數(shù)組中只需要 1 個元素，因為時間是單個浮點值 */
const timeJs = new Float32Array(1)
let t = 5.3
/* 純 JS，只需設置值 */
timeJs.set([t], 0)
/* 將數(shù)據(jù)從 CPU/JS 傳遞到 GPU/WGSL */
device.queue.writeBuffer(timeBuffer, 0, timeJs);</code>

Particle Position - WGSL Storage

We store and update particle positions directly in GPU-accessible memory – allowing us to update them in parallel by taking advantage of the GPU’s massive multi-core architecture.

Parallelization is coordinated with the help of work group size, declared in the compute 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 The value provides the thread identifier.

By definition, global_invocation_id = workgroup_id * workgroup_size local_invocation_id - this means it can be used as a particle index.

For example, if we have 10k particles and workgroup_size is 64, we need to schedule Math.ceil(10000/64) workgroups. Each time a compute pass is triggered from JS, we will explicitly tell the GPU to perform this amount of work:

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

If PARTICLE_COUNT == 10000 and WORKGROUP_SIZE == 64, we will start 157 workgroups (10000/64 = 156.25), and the calculated range of local_invocation_id of each workgroup is 0 to 63 (while the range of workgroup_id is 0 to 157 ). Since 157 * 64 = 1048, we will end up doing slightly more calculations in a workgroup. We handle overflow by discarding redundant calls.

Here is the final result of calculating the shader after taking these factors into account:

<code>@compute @workgroup_size(${WORKGROUP_SIZE})
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
  let index = global_id.x;
  // 由于工作組網(wǎng)格未對齊，因此丟棄額外的計算
  if (index >= arrayLength(&particles)) {
    return;
  }
  /* 將整數(shù)索引轉換為浮點數(shù)，以便我們可以根據(jù)索引（和時間）計算位置更新 */
  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>

These values ??will persist across calculation passes because particles are defined as storage variables.

Read the particle position in the compute shader in the vertex shader

In order to read the particle positions in the vertex shader from the compute shader, we need a read-only view, since only the compute shader can write to the storage.

The following is a statement from 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>

Trying to re-use the same read_write style in a compute shader will just error:

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

Note that the binding numbers in the vertex shader do not have to match the compute shader binding numbers - they only need to match the vertex shader's binding group declaration:

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

I selected binding:2 in the GitHub sample code http://m.miracleart.cn/link/2e5281ee978b78d6f5728aad8f28fedb#L70 - just to explore the boundaries of the constraints imposed by WebGPU

Run the simulation step by step

With all settings in place, the update and render loops are coordinated in JS:

<code>/* 從 t = 0 開始模擬 */
let t = 0
function frame() {
  /*
    為簡單起見，使用恒定整數(shù)時間步 - 無論幀速率如何，都會一致渲染。
  */
  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);
  // 重要的是要調度正確數(shù)量的工作組以處理所有粒子
  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 unleashes the power of massively parallel GPU computing in the browser.

It runs in passes - each pass has local variables enabled through a pipeline with memory binding (bridging CPU memory and GPU memory).

Compute delivery allows for the coordination of parallel workloads through workgroups.

While it does require some heavy setup, I think the local binding/state style is a huge improvement over WebGL's global state model - making it easier to use while also finally bringing the power of GPU computing to Entered the Web.

The above is the detailed content of WebGPU tutorial: compute, vertex, and fragment shaders on the web. For more information, please follow other related articles on the PHP Chinese website!