Aula 17: __syncthreads() e Condições de Corrida

Quando você carrega dados na shared memory e depois os lê entre threads diferentes, surge um problema: as threads de um block não correm em lockstep exato. Uma thread pode chegar à linha de leitura antes que outra thread termine de escrever o slot de que ela precisa — isso é uma condição de corrida

Um grupo de pessoas coloca livros numa prateleira partilhada, e depois cada um precisa pegar um livro que outra pessoa colocou. Se você começar a pegar antes que todos tenham terminado de colocar, alguns vão esticar a mão para um lugar vazio. __syncthreads() é dizer 'ninguém pega até que todos tenham colocado'.

__syncthreads(): Uma barreira dentro de um block: todas as threads param nesse ponto até que todas tenham chegado, e só então continuam juntas.
condição de corrida (race condition): Quando o resultado depende da ordem em que as threads chegam a uma escrita/leitura. Aqui: ler da shared antes que a escrita tenha terminado.
barreira: Um ponto de espera que todas as threads do block precisam alcançar antes que qualquer uma prossiga. Ela bloqueia corridas entre uma escrita e uma leitura.
escopo da barreira: __syncthreads() funciona apenas dentro de um único block. Não tem efeito entre blocks, e não deve ser chamada num branch divergente.