Qué es una storage texture y en qué se diferencia de una normal

Una storage texture es la misma GPUTexture de siempre atada de otra manera. El cambio de binding trae una capacidad —el shader puede escribir en ella— y le quita otra: no hay sampler, no hay filtrado, no hay coordenadas normalizadas, no hay selección de mip. Se accede por coordenadas enteras de texel, exactamente como a un array bidimensional, y esa es la clave de cuándo compensa usarla.

El binding

Del lado de la API, la entrada del bind group layout usa la clase storageTexture, que exige un format y admite access y viewDimension:

const layout = device.createBindGroupLayout({
  label: 'salida del compute',
  entries: [{
    binding: 0,
    visibility: GPUShaderStage.COMPUTE,
    storageTexture: {
      format: 'rgba8unorm',     // obligatorio, fijo en el layout
      access: 'write-only',     // por defecto
      viewDimension: '2d',      // por defecto
    },
  }],
});

La textura necesita GPUTextureUsage.STORAGE_BINDING en su usage, y su formato tiene que estar en la lista de formatos que admiten storage.

Del lado del shader, el tipo lleva el formato y el modo de acceso como parámetros del tipo, no como una anotación:

@group(0) @binding(0) var salida : texture_storage_2d<rgba8unorm, write>;

Fíjate en la asimetría de nombres: en la API el acceso se llama 'write-only' y en WGSL se llama write. Los tres pares son 'write-only' y write, 'read-only' y read, 'read-write' y read_write.

El recurso que se pasa en el bind group es una GPUTextureView, igual que con una textura muestreada:

const grupo = device.createBindGroup({
  layout,
  entries: [{ binding: 0, resource: destino.createView() }],
});

Lo que el binding te quita

Es más corto listar lo que pierdes que lo que ganas.

No hay sampler. Ni filtrado bilineal, ni modos de dirección, ni anisotropía. Una coordenada fuera del rango no se envuelve ni se sujeta: la escritura simplemente se descarta.

No hay coordenadas normalizadas. Se accede con enteros de texel, del 0 al ancho menos uno. Es la misma convención que textureLoad sobre una textura normal.

No hay selección de nivel de mip. La vista atada expone exactamente un nivel, y textureStore no recibe parámetro de nivel. Para escribir a varios niveles hacen falta varias vistas y varios bindings o varios dispatches.

No es visible desde la etapa de vértices. Igual que con los storage buffers de escritura, la validación rechaza un layout con storageTexture cuya visibility incluya GPUShaderStage.VERTEX. Se puede usar desde FRAGMENT y desde COMPUTE.

No admite viewDimension: 'cube' ni 'cube-array'. Para escribir a un cubemap hay que verlo como '2d-array' de seis capas.

El formato está fijado en el layout. Un layout con format: 'rgba8unorm' solo acepta vistas de ese formato exacto. No hay conversión: los bits que escribes son los bits que quedan, interpretados según el formato declarado.

El montaje mínimo

Un compute shader que genera una imagen y un render pass que la muestrea. Es el patrón base de todo el post-proceso por compute.

const ANCHO = 512, ALTO = 512;

const destino = device.createTexture({
  label: 'generada',
  size: [ANCHO, ALTO],
  format: 'rgba8unorm',
  usage: GPUTextureUsage.STORAGE_BINDING   // escribir desde el compute
       | GPUTextureUsage.TEXTURE_BINDING,  // muestrear después
});

Las dos capacidades a la vez son legítimas y frecuentes: escribir con un binding y leer con otro, en passes distintos.

@group(0) @binding(0) var salida : texture_storage_2d<rgba8unorm, write>;

struct Params { tiempo : f32 };
@group(0) @binding(1) var<uniform> params : Params;

@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) id : vec3<u32>) {
  let tam = textureDimensions(salida);
  if (id.x >= tam.x || id.y >= tam.y) { return; }

  let uv = vec2<f32>(id.xy) / vec2<f32>(tam);
  let color = vec3<f32>(
    0.5 + 0.5 * sin(uv.x * 12.0 + params.tiempo),
    0.5 + 0.5 * sin(uv.y * 12.0 + params.tiempo * 1.3),
    0.5 + 0.5 * sin((uv.x + uv.y) * 8.0 - params.tiempo));

  textureStore(salida, vec2<i32>(id.xy), vec4<f32>(color, 1.0));
}

Dos detalles importan. textureDimensions() sobre una storage texture no lleva parámetro de nivel, y devuelve un vec2<u32> para las 2D. Y la guarda de límites es obligatoria: el dispatch se redondea hacia arriba, así que sobran invocaciones en los bordes cuando las dimensiones no son múltiplo del tamaño de workgroup.

const pass = encoder.beginComputePass();
pass.setPipeline(pipelineGenerar);
pass.setBindGroup(0, grupoGenerar);
pass.dispatchWorkgroups(Math.ceil(ANCHO / 8), Math.ceil(ALTO / 8));
pass.end();

// En el mismo encoder: el render pass ve las escrituras del compute.
const rp = encoder.beginRenderPass(descriptorPass);
rp.setPipeline(pipelineDibujo);
rp.setBindGroup(0, grupoMuestreo);   // aquí 'destino' va como texture, no storage
rp.draw(3);
rp.end();

El mismo command encoder garantiza el orden y la visibilidad de las escrituras entre passes, sin barreras explícitas. La implementación inserta lo que haga falta.

El tamaño de workgroup

@workgroup_size(8, 8) no es un número mágico pero tampoco es arbitrario. La razón de que sea bidimensional y cuadrado es la localidad: las 64 invocaciones de ese grupo tocan un bloque de 8×8 texels, que en una textura con disposición en mosaico caen en unas pocas líneas de caché. Un @workgroup_size(64, 1) haría el mismo número de invocaciones tocando una tira de 64 texels de ancho, que atraviesa varios mosaicos.

Los límites garantizados son 256 invocaciones por workgroup, con máximos de 256 en X, 256 en Y y 64 en Z. Los tamaños que se usan en la práctica para trabajo sobre imágenes son 8×8 y 16×16; para trabajo unidimensional, 64 o 256.