textureStore y las funciones de acceso directo

textureStore es la única forma de escribir a una textura desde un shader, y su firma es más simple de lo que sugiere el nombre: una textura, unas coordenadas enteras y un vec4. Lo que no es simple es qué pasa entre ese vec4 y los bits que acaban en memoria, porque la conversión la decide el formato declarado en el tipo y no siempre es la que esperas.

Las sobrecargas

textureStore tiene una variante por cada dimensión de storage texture. Las coordenadas son enteras, con i32 o u32, y el valor es siempre un vec4 del tipo del canal.

// 1D: coordenada escalar.
@group(0) @binding(0) var linea : texture_storage_1d<rgba8unorm, write>;
textureStore(linea, i, vec4<f32>(c, 1.0));

// 2D: coordenada vec2.
@group(0) @binding(1) var imagen : texture_storage_2d<rgba8unorm, write>;
textureStore(imagen, vec2<i32>(x, y), vec4<f32>(c, 1.0));

// 2D array: coordenada vec2 más índice de capa como argumento aparte.
@group(0) @binding(2) var capas : texture_storage_2d_array<rgba8unorm, write>;
textureStore(capas, vec2<i32>(x, y), capa, vec4<f32>(c, 1.0));

// 3D: coordenada vec3.
@group(0) @binding(3) var volumen : texture_storage_3d<rgba16float, write>;
textureStore(volumen, vec3<i32>(x, y, z), vec4<f32>(c, 1.0));

El tipo del vec4 depende del formato. Para formatos unorm, snorm y float, es vec4<f32>. Para formatos uint, es vec4<u32>. Para formatos sint, es vec4<i32>. Pasar el tipo equivocado es un error de compilación, no una conversión implícita.

Las coordenadas fuera de límites no producen error ni corrupción: la escritura se descarta. Eso lo garantiza el modelo de comportamiento acotado de WGSL, y significa que la guarda del dispatch protege tus datos pero no la memoria del proceso, que ya está protegida.

La conversión del valor

Aquí es donde hay que tener cuidado. El vec4<f32> que escribes no se copia tal cual: se convierte al formato de destino según sus reglas.

Con un formato unorm como rgba8unorm, cada componente se sujeta al rango [0, 1] y se escala a [0, 255]. Un valor de 1,5 se guarda como 255; uno de -0,3 como 0. No hay aviso.

Con un formato snorm como rgba8snorm, el rango es [-1, 1] y se mapea a [-127, 127].

Con un formato float como rgba16float, se convierte a media precisión con el redondeo correspondiente. Los valores que exceden el rango de f16 —unos 65 504— se convierten a infinito.

Con un formato de menos de cuatro canales, los componentes sobrantes se descartan. Escribir vec4<f32>(r, g, b, a) a una texture_storage_2d<r32float, write> guarda solo r; los otros tres se ignoran en silencio. Es el comportamiento correcto según la especificación y una fuente perfecta de confusión, porque el código compila y parece correcto.

// r32float: solo se guarda la componente x. Las otras tres son decorativas.
@group(0) @binding(0) var profundidad : texture_storage_2d<r32float, write>;
textureStore(profundidad, coord, vec4<f32>(z, 0.0, 0.0, 0.0));

Leer y consultar

Leer desde una storage texture

Con acceso read o read_write, textureLoad funciona sobre storage textures y devuelve el valor sin filtrar ni convertir coordenadas:

@group(0) @binding(0) var origen : texture_storage_2d<rgba8unorm, read>;

let texel = textureLoad(origen, vec2<i32>(x, y));   // sin parámetro de nivel

Fíjate en la diferencia con textureLoad sobre una textura normal: allí la firma incluye un parámetro de nivel de mip, aquí no, porque la vista de una storage texture expone exactamente un nivel.

Estos dos modos de acceso requieren la extensión de lenguaje readonly_and_readwrite_storage_textures, disponible en navigator.gpu.wgslLanguageFeatures, y su disponibilidad es el tema de la lección sobre los modos de acceso.

Con acceso write a secas, leer es un error de compilación. Si necesitas leer y escribir la misma imagen, o usas read_write donde esté disponible, o usas dos texturas.

Las funciones de consulta

textureDimensions sobre una storage texture no lleva parámetro de nivel y devuelve un entero sin signo por dimensión:

let tam2d = textureDimensions(imagen);    // vec2<u32>
let tam3d = textureDimensions(volumen);   // vec3<u32>
let tam1d = textureDimensions(linea);     // u32

textureNumLayers devuelve el número de capas de una texture_storage_2d_array:

let capas = textureNumLayers(arrayStorage);   // u32

No existe textureNumLevels para storage textures, porque siempre tienen un nivel visible.

El patrón idiomático para la guarda de límites usa la comparación por componentes:

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

any() sobre un vec2<bool> es más compacto que dos comparaciones y compila a lo mismo.

Un ejemplo completo: desenfoque separable

El desenfoque gaussiano separable es el ejemplo canónico de storage textures porque muestra las tres cosas a la vez: lectura muestreada, escritura directa, y dos pasadas encadenadas.

@group(0) @binding(0) var origen  : texture_2d<f32>;
@group(0) @binding(1) var smp     : sampler;
@group(0) @binding(2) var destino : texture_storage_2d<rgba16float, write>;

struct Params { direccion : vec2<f32>, radio : f32 };
@group(0) @binding(3) var<uniform> params : Params;

const PESOS = array<f32, 5>(0.227027, 0.194594, 0.121621, 0.054054, 0.016216);

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

  let uv = (vec2<f32>(id.xy) + vec2<f32>(0.5)) / vec2<f32>(tam);
  let paso = params.direccion * params.radio / vec2<f32>(tam);

  var suma = textureSampleLevel(origen, smp, uv, 0.0) * PESOS[0];
  for (var i = 1; i < 5; i = i + 1) {
    let d = paso * f32(i);
    suma = suma + textureSampleLevel(origen, smp, uv + d, 0.0) * PESOS[i];
    suma = suma + textureSampleLevel(origen, smp, uv - d, 0.0) * PESOS[i];
  }

  textureStore(destino, vec2<i32>(id.xy), suma);
}

Se despacha dos veces con direccion en (1, 0) y después en (0, 1), alternando origen y destino. La lectura usa textureSampleLevel —no textureSample— porque en un compute shader no hay derivadas y por tanto no hay cálculo automático de LOD; el nivel hay que darlo explícitamente. Es un error que el compilador atrapa, afortunadamente.

Y usar el sampler bilineal para leer, en vez de textureLoad, no es capricho: permite el truco clásico de tomar muestras entre texels y aprovechar el filtro del hardware para promediar dos texels con una sola muestra, reduciendo a la mitad el número de accesos de un kernel gaussiano.