Sincronizar con __syncthreads

Los hilos de un bloque no avanzan al unísono. Uno puede terminar de escribir mientras otro todavía está cargando, así que no puede suponer que los datos están listos.

Una barrera es un punto del código en el que todos los hilos deben esperar hasta que lleguen todos. Solo entonces el bloque continúa conjuntamente.

CUDA le proporciona __syncthreads() como barrera para todo el bloque. Llámela y todos los hilos del bloque se detendrán hasta que llegue el último. 🛑

El patrón fundamental es sencillo: cada hilo carga su parte en la memoria compartida, se sincronizan y después todos leen de forma segura lo que han escrito los demás.

Si olvida la barrera, un hilo puede leer una posición antes de que su vecino haya escrito en ella. Es una condición de carrera y produce datos erróneos de forma silenciosa.

La barrera también hace que las escrituras compartidas sean visibles para todos los hilos después de ella. Por tanto, una vez superado __syncthreads, todos ven los valores recién escritos.

La regla estricta es que todos los hilos del bloque deben llegar al mismo __syncthreads. Si algunos lo omiten, el bloque puede quedar bloqueado para siempre.

Nunca coloque una barrera dentro de un if al que solo entren algunos hilos. Los hilos que sigan la otra ruta nunca llegarán y el bloque quedará bloqueado. ⚠️

Un límite importante: __syncthreads solo sincroniza un bloque. No puede coordinar hilos de bloques distintos del grid.

En los algoritmos con tiles, a menudo se sincroniza dos veces por fase: una después de cargar un tile y otra después de calcular, antes de cargar el siguiente.

Una barrera cuesta algo de tiempo mientras los hilos esperan. Úsela cuando la corrección lo exija, pero evite llamadas adicionales que solo detengan a los hilos rápidos.

Ha aprendido que __syncthreads() es una barrera para todo el bloque que mantiene sincronizados los hilos, y que colocarla dentro de ramas divergentes puede provocar un bloqueo. A continuación: conflictos de bancos. 🎯