write, read y read_write: los tres modos de acceso

Una storage texture solo escribible es un destino y nada más. Los otros dos modos —leer, y leer y escribir a la vez— llegaron después al lenguaje, requieren comprobar una extensión, y en el caso de read_write están limitados a un puñado de formatos. Saber exactamente qué está disponible en cada nivel evita construir un algoritmo sobre una capacidad que la mitad de los dispositivos no tienen.

Los tres modos y su disponibilidad

El modo por defecto en la API es 'write-only', así que omitir access en el layout da el modo básico. En WGSL no hay valor por defecto: el parámetro del tipo es obligatorio.

'read-only' y 'read-write' requieren que la extensión de lenguaje readonly_and_readwrite_storage_textures esté presente. Se comprueba antes de compilar el shader:

const tieneRW = navigator.gpu.wgslLanguageFeatures.has(
  'readonly_and_readwrite_storage_textures');

wgslLanguageFeatures es un WGSLLanguageFeatures —un set de solo lectura— que cuelga del objeto navigator.gpu, no del adaptador ni del dispositivo. Son capacidades del compilador de WGSL del navegador, no del hardware, y por eso se consultan sin haber pedido nada.

Si la extensión no está y creas un layout con access: 'read-only' o 'read-write', obtienes un GPUValidationError.

Los formatos de read_write

Aquí está la restricción que más condiciona el diseño. En el núcleo de WebGPU, el acceso read-write está disponible únicamente para los tres formatos de un canal de 32 bits: r32uint, r32sint y r32float.

Con la feature texture-formats-tier2, la lista se amplía a r8unorm, r8uint, r8sint, rgba8unorm, rgba8uint, rgba8sint, r16uint, r16sint, r16float, rgba16uint, rgba16sint, rgba16float, rgba32uint, rgba32sint y rgba32float.

El acceso read-only, en cambio, funciona con cualquier formato que admita STORAGE_BINDING, con una excepción explícita: bgra8unorm no se puede usar como storage texture de solo lectura. La especificación lo prohíbe porque ese formato existe para escritura al canvas y su lectura no es portable.

const RW_NUCLEO = new Set(['r32uint', 'r32sint', 'r32float']);
const RW_TIER2 = new Set([
  'r8unorm', 'r8uint', 'r8sint',
  'rgba8unorm', 'rgba8uint', 'rgba8sint',
  'r16uint', 'r16sint', 'r16float',
  'rgba16uint', 'rgba16sint', 'rgba16float',
  'rgba32uint', 'rgba32sint', 'rgba32float',
]);

function admiteReadWrite(device, formato) {
  if (!navigator.gpu.wgslLanguageFeatures.has(
        'readonly_and_readwrite_storage_textures')) return false;
  if (RW_NUCLEO.has(formato)) return true;
  return RW_TIER2.has(formato) && device.features.has('texture-formats-tier2');
}

Leer y escribir la misma textura

read_write permite el patrón in-place: leer un texel, modificarlo y escribirlo en la misma invocación.

// No hace falta ninguna directiva enable: la extensión se comprueba en JS.
@group(0) @binding(0) var campo : texture_storage_2d<r32float, read_write>;

@compute @workgroup_size(8, 8)
fn decaer(@builtin(global_invocation_id) id : vec3<u32>) {
  let tam = textureDimensions(campo);
  if (any(id.xy >= tam)) { return; }

  let c = vec2<i32>(id.xy);
  let v = textureLoad(campo, c).r;
  textureStore(campo, c, vec4<f32>(v * 0.98, 0.0, 0.0, 0.0));
}

Este caso es seguro porque cada invocación toca únicamente su propio texel. En cuanto una invocación lee texels vecinos que otra invocación escribe, aparece la misma clase de carrera que en los storage buffers, con las mismas consecuencias: resultado dependiente del orden de ejecución de los workgroups, que no está definido.

// INCORRECTO: lee vecinos que otras invocaciones están escribiendo.
@compute @workgroup_size(8, 8)
fn difundir(@builtin(global_invocation_id) id : vec3<u32>) {
  let c = vec2<i32>(id.xy);
  let suma = textureLoad(campo, c + vec2<i32>(-1, 0)).r
           + textureLoad(campo, c + vec2<i32>( 1, 0)).r
           + textureLoad(campo, c + vec2<i32>(0, -1)).r
           + textureLoad(campo, c + vec2<i32>(0,  1)).r;
  textureStore(campo, c, vec4<f32>(suma * 0.25, 0.0, 0.0, 0.0));
}

La solución es la misma que con buffers: dos texturas que alternan, o un read de origen y un write de destino en bindings distintos. Con la ventaja añadida de que el binding de solo lectura puede ser un texture_2d normal, y entonces puedes usar un sampler y aprovechar el filtro del hardware, cosa que la storage texture de lectura no permite.

storageBarrier() sincroniza las escrituras a var<storage> dentro de un workgroup, pero no garantiza nada entre workgroups distintos ni cubre las escrituras a storage textures del mismo modo. Para el caso de vecinos que cruzan la frontera del workgroup, no hay barrera que valga: hace falta terminar el dispatch.

Cuándo compensa read_write

El modo in-place existe porque ahorra memoria, no porque sea más rápido. Con read_write una simulación necesita una textura en vez de dos, y para un volumen de 256 al cubo en r32float eso son 64 MiB en vez de 128.

Los casos donde compensa son aquellos en que la operación es puramente local: cada texel depende solo de sí mismo. Decaimiento exponencial, aplicación de una curva, cuantización, sujetado de rango, acumulación de un valor calculado en otro sitio. En todos ellos, read_write evita la copia y evita el segundo recurso.

Los casos donde no compensa son los que leen vecinos, y son la mayoría de los algoritmos interesantes: convoluciones, difusión, propagación, gradientes. Para todos ellos, el doble buffer es obligatorio y el modo read_write no ayuda.

Hay un tercer patrón que sí saca partido: usar read_write como acumulador donde cada invocación solo escribe su propio texel pero necesita leer lo que hay para acumular. Un mapa de calor al que varios dispatches suman contribuciones, por ejemplo. Y ahí conviene recordar que si dos invocaciones del mismo dispatch pudieran escribir el mismo texel, hace falta atomicidad, y las storage textures no ofrecen operaciones atómicas en WebGPU.