wandres.dev
STORAGE BUFFERS · Lectura y escritura arbitraria

Arrays de tamaño en tiempo de ejecución y arrayLength()

El único tipo de WGSL cuyo tamaño no se conoce al compilar: dónde puede aparecer, cómo se calcula su longitud, y el patrón de cabecera y cuerpo que resuelve el caso general.

⏱ 18 min

Todos los tipos de WGSL tienen tamaño conocido en tiempo de compilación menos uno: el array sin longitud. Es la característica que permite que un mismo shader procese diez partículas o un millón sin recompilar, y viene con tres restricciones muy concretas sobre dónde puede aparecer. Entenderlas evita el rodeo de pasar la longitud en un uniform aparte, que es lo que hace todo el mundo antes de descubrir arrayLength().

🎯 Al terminar esta lección sabrás
  • Declarar un array de tamaño en tiempo de ejecución y conocer sus tres restricciones.
  • Calcular su longitud en el shader con arrayLength() y su forma de puntero.
  • Explicar de dónde sale ese número y por qué depende del size del binding.
  • Aplicar el patrón de cabecera y cuerpo cuando hacen falta metadatos junto al array.

Dónde puede aparecer

Un array de tamaño en tiempo de ejecución se escribe sin el segundo parámetro:

@group(0) @binding(0) var<storage, read> particulas : array<Particula>;

Las tres restricciones son estas.

Solo en el espacio storage. Ni uniform, ni workgroup, ni private, ni function. El espacio uniform exige tamaño conocido porque el hardware lo coloca en la caché de constantes, cuyo tamaño se reserva al crear el pipeline.

Solo como variable de módulo, o como último miembro de una struct que sea el tipo de una variable de módulo. No puede haber nada detrás, porque nada tendría offset conocido.

Solo uno por struct. Consecuencia de lo anterior: si es el último miembro, solo puede haber uno.

Este es el caso válido con struct:

struct Sistema {
  gravedad  : vec3<f32>,
  dt        : f32,
  particulas: array<Particula>,   // último miembro, sin tamaño
};
@group(0) @binding(0) var<storage, read_write> sistema : Sistema;

Y estos dos son errores de compilación:

struct Mal1 {
  datos : array<f32>,   // no es el último
  n     : u32,
};

struct Mal2 {
  a : array<f32>,       // dos arrays sin tamaño
  b : array<f32>,
};

arrayLength y el puntero

arrayLength() devuelve el número de elementos del array, y su argumento es un puntero al array, no el array:

let n = arrayLength(&particulas);          // variable de módulo
let m = arrayLength(&sistema.particulas);  // miembro de struct

El operador & es obligatorio. Pasar el array directamente es un error de tipo, y es el fallo más frecuente al usar esta función por primera vez. El resultado es un u32.

El uso canónico es la guarda del compute shader. Como los workgroups se despachan en múltiplos del tamaño del grupo, casi siempre sobran invocaciones al final, y sin la guarda escriben fuera del array:

@compute @workgroup_size(64)
fn integrar(@builtin(global_invocation_id) id : vec3<u32>) {
  let i = id.x;
  if (i >= arrayLength(&particulas)) { return; }
  particulas[i].pos = particulas[i].pos + particulas[i].vel * sistema.dt;
}

Y del lado de JavaScript, el número de workgroups se redondea hacia arriba:

const TAM_GRUPO = 64;
pass.dispatchWorkgroups(Math.ceil(numParticulas / TAM_GRUPO));
💡
La guarda no es opcional aunque el acceso esté acotado

WGSL garantiza que un acceso fuera de límites a un storage buffer no lee ni escribe memoria ajena: la implementación lo redirige dentro del rango o lo descarta. Pero redirigido dentro del rango significa que puede escribir sobre un elemento válido, corrompiendo datos reales. La guarda no protege la memoria del proceso, que ya está protegida: protege tus propios datos.

De dónde sale ese número

arrayLength() no lee ninguna variable que tú hayas escrito. La implementación calcula el valor a partir del tamaño real del binding en el momento del draw o del dispatch, con esta fórmula:

arrayLength = floor((tamañoDelBinding - offsetDelArrayDentroDeLaStruct) / strideDelElemento)

El tamañoDelBinding es el size que pasaste en el GPUBufferBinding del bind group, o lo que quede del buffer desde offset si omitiste size. Eso tiene una consecuencia práctica muy útil: puedes cambiar la longitud vista por el shader sin tocar el buffer ni el shader, solo creando un bind group con otro size.

// El mismo buffer, dos vistas de longitudes distintas.
const grupoTodas = device.createBindGroup({
  layout, entries: [{ binding: 0, resource: { buffer: bufParticulas } }],
});
const grupoMitad = device.createBindGroup({
  layout, entries: [{ binding: 0,
    resource: { buffer: bufParticulas, offset: 0, size: bufParticulas.size / 2 } }],
});

Y una consecuencia menos agradable: si el tamaño del binding no es múltiplo exacto del stride, la división trunca y sobra memoria al final. Con structs cuyo tamaño no divide limpiamente el tamaño del buffer, la última entrada puede quedar fuera del recuento. Dimensionar el buffer como stride × n exacto elimina el problema.

El patrón de cabecera y cuerpo

Cuando hacen falta metadatos junto al array —un contador de elementos vivos, parámetros de la simulación—, hay dos formas y conviene elegir con criterio.

La primera es meterlo todo en una struct con el array al final, como en el ejemplo de Sistema de arriba. Es compacta y evita un binding, pero mezcla un dato que cambia por frame con un array enorme que casi no cambia, y obliga a escribir en el mismo buffer los dos.

La segunda es separar: un uniform pequeño con la cabecera y un storage con el cuerpo.

struct Params { gravedad : vec3<f32>, dt : f32, vivas : u32 };

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

Esta versión gasta un binding más y a cambio permite escribir los 32 bytes de parámetros cada frame sin tocar los 32 MiB de partículas, deja la cabecera en el camino de constantes con difusión, y hace que arrayLength() refleje la capacidad del buffer mientras params.vivas refleja cuántas están en uso. Esa distinción entre capacidad y ocupación es la que aparece en cuanto la simulación crea y destruye elementos.

@compute @workgroup_size(64)
fn integrar(@builtin(global_invocation_id) id : vec3<u32>) {
  let i = id.x;
  if (i >= params.vivas) { return; }        // ocupación, no capacidad
  // ...
}
El contador tiene que vivir en la GPU si la GPU lo modifica

El patrón de arriba funciona mientras sea la CPU quien decide cuántas partículas hay. En cuanto la propia simulación las crea o las destruye —un emisor que añade partículas, un culling que descarta objetos—, el contador tiene que ser un atomic<u32> en un storage buffer, porque quien lo incrementa es el shader. Y entonces aparece el problema de verdad: la CPU ya no sabe cuántos elementos hay, así que no puede calcular cuántos workgroups despachar. La salida no es leer el contador de vuelta a la CPU, que cuesta un frame entero de latencia, sino dispatchWorkgroupsIndirect(): un buffer con el número de workgroups que otro compute shader escribió a partir del contador atómico. Ese es el momento en que el pipeline deja de tener a la CPU en el camino crítico, y todo empieza con la decisión de dónde vive un u32.

⚔️ Reto práctico

Crea un storage buffer de 1024 structs y ata tres bind groups distintos sobre él con size de 1024, 512 y 100 veces el stride. Imprime arrayLength() desde un compute shader en los tres casos y comprueba la fórmula de truncamiento con un tamaño que no sea múltiplo exacto del stride.