El patrón de doble buffer para simulaciones

En cuanto una simulación necesita leer el estado de los vecinos para calcular el suyo, escribir sobre el mismo buffer deja de funcionar: no hay forma de saber si el vecino que lees es del paso anterior o del actual. La solución es tan vieja como la computación gráfica y sigue siendo la correcta: dos buffers, uno de lectura y otro de escritura, que intercambian su papel en cada paso. Lo interesante en WebGPU es que el intercambio no toca ningún recurso, solo el bind group que se ata.

Cuándo hace falta y cuándo no

La pregunta es una sola: ¿una invocación lee alguna posición que otra invocación escribe en el mismo dispatch?

Si cada invocación solo toca su propio índice, no hace falta doble buffer. Una integración de partículas independientes —posición más velocidad por delta de tiempo— es exactamente ese caso, y puede escribirse en su sitio con read_write sobre un solo buffer.

Si la invocación lee vecinos, hace falta. Un autómata celular lee las ocho celdas adyacentes; un desenfoque separable lee un radio a cada lado; una simulación de fluidos SPH lee todas las partículas dentro de un radio; un paso de relajación de restricciones lee las dos partículas de cada enlace. En todos ellos, sin doble buffer, el resultado depende del orden en que la GPU decidió ejecutar los workgroups, que no está definido y cambia entre ejecuciones.

Los dos buffers y los dos grupos

El montaje es mecánico. Dos buffers idénticos, y dos bind groups que los enchufan al revés:

const TAM = numParticulas * STRIDE;

const buffers = [0, 1].map((i) => device.createBuffer({
  label: `estado ${i}`,
  size: TAM,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.VERTEX,
}));

device.queue.writeBuffer(buffers[0], 0, estadoInicial);

// El grupo 0 lee de buffers[0] y escribe en buffers[1]; el grupo 1 al revés.
const grupos = [0, 1].map((i) => device.createBindGroup({
  label: `sim ${i}`,
  layout: layoutSim,
  entries: [
    { binding: 0, resource: { buffer: bufferParams } },
    { binding: 1, resource: { buffer: buffers[i] } },       // lectura
    { binding: 2, resource: { buffer: buffers[1 - i] } },   // escritura
  ],
}));

El layout declara los dos storage con modos distintos:

const layoutSim = device.createBindGroupLayout({
  label: 'simulacion',
  entries: [
    { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
    { binding: 1, visibility: GPUShaderStage.COMPUTE,
      buffer: { type: 'read-only-storage' } },
    { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
  ],
});

Y el shader, que nunca se entera del intercambio:

struct Particula { pos : vec2<f32>, vel : vec2<f32> };

@group(0) @binding(0) var<uniform>              params  : Params;
@group(0) @binding(1) var<storage, read>        entrada : array<Particula>;
@group(0) @binding(2) var<storage, read_write>  salida  : array<Particula>;

@compute @workgroup_size(64)
fn paso(@builtin(global_invocation_id) id : vec3<u32>) {
  let i = id.x;
  let n = arrayLength(&entrada);
  if (i >= n) { return; }

  var p = entrada[i];
  var fuerza = vec2<f32>(0.0, -9.8);

  // Lee vecinos del buffer de ENTRADA: consistente, del paso anterior.
  for (var j = 0u; j < n; j = j + 1u) {
    if (j == i) { continue; }
    let d = entrada[j].pos - p.pos;
    let r2 = max(dot(d, d), 0.0001);
    fuerza = fuerza - normalize(d) * params.repulsion / r2;
  }

  p.vel = p.vel + fuerza * params.dt;
  p.pos = p.pos + p.vel * params.dt;
  salida[i] = p;
}

La separación entre entrada y salida es lo que hace correcta la lectura de vecinos: todas las invocaciones ven el mismo estado congelado del paso anterior.

El bucle de simulación

El intercambio es una variable entera. No se recrea nada, no se copia nada:

let actual = 0;

function frame() {
  device.queue.writeBuffer(bufferParams, 0, params);

  const encoder = device.createCommandEncoder();

  const cp = encoder.beginComputePass();
  cp.setPipeline(pipelineSim);
  cp.setBindGroup(0, grupos[actual]);
  cp.dispatchWorkgroups(Math.ceil(numParticulas / 64));
  cp.end();

  // Tras el dispatch, el estado nuevo está en buffers[1 - actual].
  const rp = encoder.beginRenderPass(descriptorPass);
  rp.setPipeline(pipelineDibujo);
  rp.setVertexBuffer(0, buffers[1 - actual]);   // el recién escrito
  rp.draw(4, numParticulas);
  rp.end();

  device.queue.submit([encoder.finish()]);
  actual = 1 - actual;
  requestAnimationFrame(frame);
}

Dos detalles importan. El primero es que el compute pass y el render pass van en el mismo command encoder: WebGPU garantiza que los passes de un command buffer se ejecutan en orden y que las escrituras de uno son visibles para el siguiente, sin barreras explícitas. La sincronización entre passes la resuelve la implementación.

El segundo es el flag GPUBufferUsage.VERTEX en los buffers de estado. Permite atar el mismo buffer que el compute acaba de escribir como buffer de vértices para el dibujado, sin ninguna copia. Es el patrón que hace que un sistema de partículas viva íntegramente en la GPU. La alternativa es atarlo como read-only-storage en el vertex shader e indexarlo con @builtin(instance_index), que evita el flag VERTEX a cambio de un binding.

El coste y cómo esquivarlo a veces

El doble buffer duplica la memoria del estado. Con un millón de partículas de 32 bytes son 32 MiB que pasan a 64. En una GPU integrada con memoria compartida, eso puede ser la diferencia entre caber y no caber.

Hay tres formas de reducirlo cuando aprieta.

Separar los campos que necesitan doble buffer de los que no. Si la posición se lee de los vecinos pero el color no, solo la posición necesita duplicarse. Partir la struct en dos buffers —uno duplicado y otro no— ahorra según la proporción.

Reducir la precisión del buffer duplicado. Posiciones en f16 con la feature shader-f16, o en enteros de punto fijo de 16 bits normalizados sobre el dominio conocido de la simulación. La mitad de memoria y, a menudo, precisión suficiente.

Comprobar si la lectura de vecinos es realmente necesaria. Muchas simulaciones leen vecinos por comodidad y podrían reestructurarse para que cada invocación acumule sobre sí misma. Un sistema de muelles, por ejemplo, puede procesarse por enlaces con atómicos en vez de por partículas con lectura de vecinos.