Caçando condições de corrida com racecheck

Uma condição de corrida de dados ocorre quando duas threads acessam a mesma memória compartilhada e pelo menos uma delas escreve nela, sem uma ordenação entre as operações. 😬

As condições de corrida frequentemente produzem a resposta correta nos testes e a resposta errada em produção, porque o resultado depende do momento imprevisível das threads.

A ferramenta racecheck do compute-sanitizer analisa os acessos à memória compartilhada e informa riscos que uma execução normal ocultaria.

Escrever em um vetor compartilhado e depois ler a posição de um vizinho sem uma barreira cria uma condição de corrida, pois a posição do vizinho talvez ainda não tenha sido escrita.

Insira __syncthreads entre a escrita e a leitura para que todas as threads do bloco terminem de escrever antes que qualquer uma leia.

O racecheck também detecta um risco de escrita após escrita, quando duas threads armazenam dados na mesma posição compartilhada sem uma ordem entre elas.

O risco mais comum é a leitura após escrita: uma thread lê um valor que outra ainda está escrevendo e, por isso, lê dados obsoletos.

O racecheck informa o tipo de risco, além das duas threads e do endereço compartilhado envolvidos, para que você saiba exatamente qual barreira está faltando.

O racecheck concentra-se na memória __shared__ dentro de um bloco; as condições de corrida globais geralmente exigem lógica cuidadosa ou operações atômicas.

Assim como no memcheck, compilar com -lineinfo permite que o racecheck aponte para a linha exata do código-fonte de cada risco.

Execute o racecheck novamente depois de adicionar a barreira. Um relatório com zero riscos significa que o acesso à memória compartilhada agora está devidamente ordenado.

Você usou o racecheck para encontrar riscos na memória compartilhada, leu os relatórios e os corrigiu com barreiras __syncthreads. Bugs de temporização, revelados. 🎯