Synchroniser avec __syncthreads

Threads in a block do not march in lockstep. One may finish writing while another is still loading, so you cannot assume the data is ready yet.

A barrier is a line in the code where every thread must wait until all of them arrive. Only then does the block move on together.

CUDA gives you __syncthreads() as the block-wide barrier. Call it and every thread in the block pauses until the last one shows up. 🛑

The golden pattern is simple: every thread loads its piece into shared memory, you sync, then everyone safely reads what others wrote.

Forget the barrier and a thread may read a slot before its neighbor wrote it. That is a race condition, and the result is silent, wrong data.

The barrier also makes shared writes visible to all threads after it. So once you pass __syncthreads, everyone sees the freshly written values.

The strict rule: every thread in the block must hit the same __syncthreads. If some skip it, the block can hang forever.

Never put a barrier inside an if that only some threads enter. Threads taking the other path never arrive, and the block deadlocks. ⚠️

One important limit: __syncthreads only synchronizes one block. It cannot coordinate threads across different blocks of the grid.

In tiled algorithms you often sync twice per phase: once after loading a tile and once after computing, before loading the next.

A barrier costs a little time while threads wait. Use it where correctness needs it, but avoid extra calls that just stall fast threads.

Let us test your grasp of the barrier.

You learned that __syncthreads() is a block-wide barrier that keeps threads in step, and that placing it inside divergent branches can deadlock. Next: bank conflicts. 🎯