workgroupBarrier y storageBarrier: qué garantiza cada una

Una barrera hace dos cosas a la vez y casi todo el mundo solo conoce una. Sincroniza la ejecución: ninguna invocación del grupo pasa de la barrera hasta que todas han llegado. Y ordena la memoria: todo lo que se escribió antes de la barrera es visible para todas las invocaciones después de ella. Sin la primera propiedad, el algoritmo se descoordina. Sin la segunda, cada invocación ve su propia versión del pasado. Las dos funciones de WGSL que las proporcionan se diferencian exactamente en sobre qué memoria actúan.

Las dos mitades de una barrera

La mitad de ejecución es la que se entiende sola: es un punto de encuentro. Todas las invocaciones del workgroup se paran ahí hasta que llega la última, y entonces siguen todas. En el hardware esto no es tan caro como parece, porque las invocaciones de un mismo grupo SIMD ya van sincronizadas por construcción; el coste real es sincronizar los distintos grupos SIMD que forman el workgroup, y con 64 invocaciones sobre warps de 32 son solo dos.

La mitad de memoria es la que produce los bugs. Una GPU reordena accesos, mantiene escrituras en cachés y colas de escritura, y no garantiza por sí sola que lo que la invocación 5 escribió sea visible para la invocación 12. La barrera actúa como valla: obliga a completar las escrituras pendientes en el address space afectado y a invalidar las lecturas cacheadas, de modo que después de la barrera todos ven lo mismo.

Y ahí está la diferencia entre las dos funciones. Ambas hacen la mitad de ejecución de forma idéntica. Se diferencian en qué address space ordenan:

Ninguna de las tres toma argumentos ni devuelve nada. Y ninguna de las tres tiene el menor efecto fuera del workgroup: storageBarrier no sincroniza workgroups distintos, por mucho que su nombre sugiera memoria global. Ese malentendido es tan frecuente y tan grave que tiene su propia lección.

Se pueden combinar. Si un tramo de código escribe tanto en memoria compartida como en un storage buffer y quieres publicar las dos cosas, llamas a las dos:

tile[li] = entrada[gid.x];
resultados[gid.x] = calculo(li);
workgroupBarrier();
storageBarrier();
// aqui todos ven tanto tile como resultados

El requisito de control de flujo uniforme

Esta es la regla dura: una barrera debe ejecutarse en control de flujo uniforme dentro del workgroup. Uniforme significa que la decisión de llegar a esa barrera, y el número de veces que se llega, es el mismo para todas las invocaciones del grupo.

Si una parte del grupo entra en un if que contiene una barrera y otra no, las que entran se quedan esperando a compañeras que nunca van a aparecer. WGSL declara ese caso comportamiento indefinido; en la práctica se manifiesta como resultados corruptos o como un cuelgue del dispositivo que el navegador termina cortando con una pérdida de contexto.

// MAL: la barrera depende de los datos.
if (entrada[gid.x] > 0.0) {
  tile[li] = entrada[gid.x];
  workgroupBarrier();          // indefinido
}

// BIEN: la condicion decide que se escribe, no si se sincroniza.
var v = 0.0;
if (entrada[gid.x] > 0.0) { v = entrada[gid.x]; }
tile[li] = v;
workgroupBarrier();            // todas llegan siempre

Lo mismo con los bucles. La barrera puede estar dentro de un bucle, pero el número de iteraciones tiene que ser el mismo para todas las invocaciones:

// BIEN: los limites son constantes del modulo.
for (var s: u32 = TAM / 2u; s > 0u; s >>= 1u) {
  if (li < s) { parcial[li] += parcial[li + s]; }
  workgroupBarrier();          // fuera del if, dentro del for uniforme
}

// MAL: cada invocacion itera un numero distinto de veces.
for (var k: u32 = 0u; k < vecinos[gid.x]; k = k + 1u) {
  acumular(k);
  workgroupBarrier();          // indefinido
}

Y el corolario que ya apareció al colocar el guardia: un return temprano hace no uniforme todo lo que viene después. En un kernel con barreras no se sale antes de tiempo; se carga el elemento neutro y se participa.

Hay un caso legítimo que confunde: la barrera dentro de un if (li == 0u) está prohibida, pero la barrera antes o después de ese bloque está perfectamente bien. El patrón habitual de “una invocación hace algo especial y luego todas lo ven” se escribe así:

if (li == 0u) {
  total = suma;                // solo una escribe
}
workgroupBarrier();            // todas sincronizan
let t = total;                 // todas leen el mismo valor

workgroupUniformLoad

Existe un caso concreto en el que el patrón anterior no basta y WGSL da una función específica: cuando el valor leído de memoria compartida va a usarse como condición de control de flujo uniforme, por ejemplo como límite de un bucle que contiene otra barrera.

El compilador no puede demostrar por análisis estático que el valor leído sea igual en todas las invocaciones, aunque tú sepas que lo es. workgroupUniformLoad resuelve las dos cosas de golpe: hace la barrera y devuelve un valor que el análisis de uniformidad acepta como uniforme.

var<workgroup> cuantos: u32;

@compute @workgroup_size(64)
fn main(@builtin(local_invocation_index) li: u32) {
  if (li == 0u) { cuantos = calcularCuantos(); }

  // Barrera + lectura, y el resultado es uniforme para el analizador.
  let n = workgroupUniformLoad(&cuantos);

  for (var k: u32 = 0u; k < n; k = k + 1u) {
    trabajo(k);
    workgroupBarrier();        // legal: n es uniforme
  }
}

Sin workgroupUniformLoad, ese bucle no compilaría: el compilador rechazaría la barrera por no poder garantizar la uniformidad del límite. Es la herramienta correcta para cualquier algoritmo cuyo número de pasos lo decide el propio grupo en ejecución, como una búsqueda con terminación temprana o un recorrido de estructura de datos.

Toma un puntero al address space workgroup y devuelve el valor apuntado. No sirve para punteros a storage; para eso hay que combinar storageBarrier con una lectura y aceptar que el análisis de uniformidad no te va a dejar usar el resultado como condición de una barrera posterior.