Урок 17: __syncthreads() и состояния гонки

Когда вы размещаете данные в shared-память и затем читаете их разными потоками, возникает проблема: потоки в block'е не выполняются в точном согласованном темпе. Один поток может дойти до строки чтения раньше, чем другой поток закончил запись нужной ему ячейки, — это состояние гонки (race condition)

Группа людей кладёт книги на общую полку, а затем каждому нужно взять книгу, которую положил кто-то другой. Если начать брать до того, как все закончили класть, некоторые потянутся за пустым местом. __syncthreads() говорит «никто не берёт, пока все не положили».

__syncthreads(): Барьер внутри block'а: все потоки останавливаются в этой точке, пока не прибудут все, и только затем продолжают вместе.
состояние гонки (race condition): Когда результат зависит от порядка, в котором потоки доходят до записи/чтения. Здесь: чтение из shared до того, как запись завершилась.
барьер (barrier): Точка ожидания, до которой все потоки в block'е должны дойти, прежде чем кто-либо продолжит. Она блокирует гонки между записью и чтением.
область действия барьера (barrier scope): __syncthreads() работает только внутри одного block'а. Он не действует между block'ами и не должен вызываться в расходящейся ветви.