Generar el draw indirect desde la GPU

Todo lo anterior sirve para llegar aquí: un buffer de veinte bytes que la GPU escribe y que la GPU consume, sin que el hilo principal sepa cuántos objetos hay ni cuáles son. Los argumentos de dibujo dejan de ser parámetros de una llamada de JavaScript y pasan a ser datos en memoria de vídeo. El formato es rígido y está especificado al byte, y las limitaciones del núcleo de WebGPU marcan exactamente hasta dónde llega esta arquitectura en la web de hoy.

El formato del buffer indirecto

Un buffer indirecto es memoria plana con enteros de 32 bits en un orden fijo. WebGPU define tres formatos, uno por cada llamada:

Los valores son enteros sin signo de 32 bits, empaquetados sin huecos, y van en el mismo orden que los argumentos del método directo equivalente. Son 16, 20 y 12 bytes respectivamente.

Las reglas que hay que respetar son tres y la validación te las recuerda:

El buffer necesita GPUBufferUsage.INDIRECT entre sus usos. Como normalmente el mismo buffer también lo escribe un compute shader, lleva además STORAGE, y a menudo COPY_DST para poder ponerlo a un estado conocido desde la CPU.

El indirectOffset tiene que ser múltiplo de 4. Los bloques miden 12, 16 o 20 bytes, todos múltiplos de 4, así que colocar bloques consecutivos en un mismo buffer no da problemas.

Un firstInstance distinto de cero requiere la feature indirect-first-instance, que hay que pedir explícitamente al crear el dispositivo y que no está en todos los adaptadores. Sin ella, escribir un valor distinto de cero en ese campo es un error de validación. Más adelante verás que hay una forma limpia de no necesitarla nunca.

const args = device.createBuffer({
  label: 'args indirectos',
  size: 5 * 4,                                    // drawIndexedIndirect
  usage: GPUBufferUsage.INDIRECT
       | GPUBufferUsage.STORAGE
       | GPUBufferUsage.COPY_DST,
});

El compute que lo rellena es trivial: una sola invocación que copia el contador que dejó la compactación en el campo instanceCount.

struct ArgsIndexado {
  indexCount    : u32,
  instanceCount : u32,
  firstIndex    : u32,
  baseVertex    : u32,
  firstInstance : u32,
};

@group(0) @binding(0) var<storage, read_write> args     : ArgsIndexado;
@group(0) @binding(1) var<storage, read>       contador : u32;

@compute @workgroup_size(1)
fn preparar() {
  args.indexCount    = 36u;         // los 36 indices de un cubo
  args.instanceCount = contador;    // lo unico que decide la GPU
  args.firstIndex    = 0u;
  args.baseVertex    = 0u;
  args.firstInstance = 0u;          // cero: sin indirect-first-instance
}

Y la llamada en el render pass no sabe nada:

pase.setPipeline(pipeline);
pase.setBindGroup(0, grupoEscena);
pase.setVertexBuffer(0, vertices);
pase.setIndexBuffer(indices, 'uint32');
pase.drawIndexedIndirect(args, 0);

Fíjate en lo que no aparece: ningún número. La CPU ha grabado cinco comandos y no ha leído ni una transformada, ni un volumen envolvente, ni el resultado de un test. Si el culling decide que hay 12 instancias visibles, se dibujan 12; si decide que hay 40.000, se dibujan 40.000. El comando grabado es idéntico.

dispatchWorkgroupsIndirect cierra el círculo por el otro lado. Si un compute pass produce una cantidad variable de trabajo que consume el siguiente compute pass, el número de workgroups del segundo también sale de un buffer:

struct ArgsDispatch { x : u32, y : u32, z : u32, };

@group(0) @binding(0) var<storage, read_write> disp     : ArgsDispatch;
@group(0) @binding(1) var<storage, read>       contador : u32;

@compute @workgroup_size(1)
fn dimensionar() {
  disp.x = (contador + 63u) / 64u;   // ceil(contador / 64) con enteros
  disp.y = 1u;
  disp.z = 1u;
}

Con eso, la fase dos del culling de oclusión solo lanza workgroups para los objetos que la fase uno descartó, en vez de para los cien mil. La cadena entera —culling, compactación, dibujo— queda sin un solo número que haya tenido que cruzar a la CPU.

El patrón completo

Este es el fotograma con todas las piezas montadas.

flowchart TB
A[Buffer de instancias con transformada y volumen envolvente] --> B[Compute de frustum]
B --> C[Array de flags de visibilidad]
H[Piramide Hi-Z del fotograma anterior] --> D[Compute de oclusion]
C --> D
D --> E[Compactacion con scan o contador atomico]
E --> F[Lista densa de indices visibles]
E --> G[Compute que escribe el buffer indirecto]
G --> I[Buffer indirecto con instanceCount por lote]
F --> R[Render pass con drawIndexedIndirect]
I --> R
R --> Z[Depth buffer del fotograma]
Z --> P[Reconstruir la piramide Hi-Z]
P --> H
style A fill:#94e2d5,color:#11111b
style H fill:#94e2d5,color:#11111b
style Z fill:#94e2d5,color:#11111b
style B fill:#fab387,color:#11111b
style D fill:#fab387,color:#11111b
style E fill:#fab387,color:#11111b
style G fill:#fab387,color:#11111b
style P fill:#fab387,color:#11111b
style C fill:#89b4fa,color:#11111b
style F fill:#89b4fa,color:#11111b
style I fill:#f9e2af,color:#11111b
style R fill:#a6e3a1,color:#11111b

Los datos entran una vez por el buffer de instancias, pasan por el frustum, por la oclusión y por la compactación, y salen convertidos en argumentos de dibujo. El único bucle que cruza fotogramas es el de la pirámide, y se cierra dentro de la GPU.

Lo que graba la CPU es esto, y no cambia nunca:

const encoder = device.createCommandEncoder();

const cull = encoder.beginComputePass({ label: 'culling' });
cull.setPipeline(pFrustum);   cull.setBindGroup(0, gFrustum);
cull.dispatchWorkgroups(gruposCull);
cull.setPipeline(pOclusion);  cull.setBindGroup(0, gOclusion);
cull.dispatchWorkgroups(gruposCull);
cull.setPipeline(pScan);      cull.setBindGroup(0, gScan);
cull.dispatchWorkgroups(gruposScan);
cull.setPipeline(pArgs);      cull.setBindGroup(0, gArgs);
cull.dispatchWorkgroups(1);
cull.end();

const render = encoder.beginRenderPass(descriptorRender);
render.setPipeline(pipeline);
render.setBindGroup(0, grupoEscena);
render.setVertexBuffer(0, megaVertices);
render.setIndexBuffer(megaIndices, 'uint32');
for (let b = 0; b < numLotes; b++) {
  render.drawIndexedIndirect(args, b * 20);    // 20 bytes por bloque
}
render.end();

device.queue.submit([encoder.finish()]);

numLotes es el número de combinaciones distintas de malla y material, no el número de objetos. Y ese es el punto entero de la arquitectura: la longitud del bucle depende de la variedad del arte, no de la cantidad de escena. Un megaVertices y un megaIndices únicos —todas las mallas concatenadas en el mismo par de buffers, cada lote apuntando a su tramo con firstIndex y baseVertex— eliminan también los setVertexBuffer por lote.

Lo que WebGPU no tiene

Toca la parte honesta. El núcleo de WebGPU no incluye multi-draw indirect. No existe ningún método que consuma un array de bloques de argumentos y un contador para emitir N dibujos con una sola llamada. En Chrome hay soporte experimental detrás de un flag, y eso significa exactamente lo que parece: no es especificación, no está en los tres motores, y no puedes construir sobre ello.

La consecuencia directa: sigues necesitando una llamada de JavaScript por lote. El bucle de drawIndexedIndirect del código de arriba no se puede colapsar. Y la consecuencia de segundo orden, más sutil: el número de lotes tiene que ser conocido por la CPU, porque es el límite del bucle. La GPU puede decidir cuántas instancias tiene cada lote, incluso puede dejarlo a cero, pero no puede inventar lotes nuevos ni eliminar la llamada de uno vacío.

La forma de vivir con esto se llama agrupar por instancia. Si muchos objetos comparten geometría, no necesitas un lote por objeto: uno solo, con instancing, y un buffer de datos por instancia indexado desde el vertex shader.

struct Instancia {
  modelo   : mat4x4<f32>,
  color    : vec4<f32>,
  material : u32,
  relleno  : vec3<u32>,
};

@group(0) @binding(0) var<uniform>       camara     : Camara;
@group(0) @binding(1) var<storage, read> instancias : array<Instancia>;
@group(0) @binding(2) var<storage, read> visibles   : array<u32>;

struct Salida {
  @builtin(position) pos   : vec4<f32>,
  @location(0)       color : vec4<f32>,
};

@vertex
fn vs(
  @location(0) posLocal : vec3<f32>,
  @builtin(instance_index) ii : u32,
) -> Salida {
  // La indireccion clave: instance_index va de 0 a instanceCount-1 y la lista
  // compactada lo traduce al indice real dentro del buffer de instancias.
  let inst = instancias[visibles[ii]];

  var s : Salida;
  s.pos   = camara.viewProj * inst.modelo * vec4<f32>(posLocal, 1.0);
  s.color = inst.color;
  return s;
}

Esa línea instancias[visibles[ii]] es el corazón de todo el capítulo. instance_index es siempre denso y empieza en cero; la lista compactada lo convierte en el índice disperso original. Y como nunca necesitas desplazar el rango de instancias, firstInstance se queda en cero y la feature indirect-first-instance deja de ser un requisito. Lo que en otras APIs se resuelve con un desplazamiento de instancia, aquí se resuelve con una indirección de un acceso a memoria, y sale más portable.

Con esto, una escena de 200.000 objetos repartidos en 400 combinaciones distintas de malla y material cuesta 400 llamadas de dibujo. A 1,5 µs por llamada son 0,6 ms de CPU, fijos: el mismo coste si mañana la escena tiene 2 millones de objetos.

La arquitectura entera, y qué cabe hoy

Un motor moderno completo apila cuatro capas sobre esto, y merece la pena saber dónde está la frontera de lo que la web puede hacer hoy.

Culling jerárquico de dos niveles. Los objetos se agrupan en clústeres de unos 64 triángulos, y el culling se hace primero por objeto y después por clúster, con una jerarquía de volúmenes envolventes. La razón es que descartar un edificio entero está bien, pero la mitad del edificio que da a la calle contraria también sobra. Cabe en WebGPU: son más storage buffers y un dispatch indirecto más. Es la mejora con mejor relación entre esfuerzo y ganancia que puedes añadir a lo de esta lección.

Mesh shaders. Sustituyen la etapa de vértices por un modelo de cómputo que emite geometría directamente, y son lo que permite que el culling de clústeres alimente al rasterizador sin pasar por un buffer intermedio. No están en WebGPU ni hay una propuesta madura. Sin ellos, el culling de clústeres tiene que materializar la lista de triángulos supervivientes en un buffer y dibujarla, que funciona pero cuesta ancho de banda.

Bindless. Un único array gigante de descriptores de textura indexado desde el shader, para que un solo pipeline pueda dibujar objetos con materiales distintos sin cambiar de bind group. No está en el núcleo. Los sustitutos razonables son los atlas de textura y los texture_2d_array, que cubren bastante si controlas el pipeline de assets, y los maxStorageBuffersPerShaderStage son 8 por defecto, así que la organización de recursos hay que pensarla.

Virtualización de geometría. Elegir el nivel de detalle por clúster en función de su tamaño en pantalla, con transiciones sin costuras. Necesita mesh shaders, o un rasterizador por software en compute, y una infraestructura de streaming considerable. Fuera de alcance en la web hoy, y probablemente durante bastante tiempo.

Lo que sí cabe entero es la columna vertebral: escena en storage buffers, culling por frustum y por oclusión en compute, compactación, argumentos indirectos generados por la GPU, y un bucle de dibujo cuya longitud depende del arte y no de la escena. Eso es el 90 % del beneficio, y está disponible en los tres motores desde que Safari 26 implementó WebGPU en septiembre de 2025. La reserva sigue siendo de cobertura, no de capacidad: WebGL2 llega a más dispositivos, así que un producto serio mantiene el camino de WebGL2 como base y trata esto como mejora progresiva.