Поиск состояний гонки с помощью racecheck

Гонка данных возникает, когда два потока обращаются к одной общей памяти и хотя бы один из них выполняет запись, а порядок этих обращений не задан. 😬

При тестировании гонки часто дают правильный результат, а в рабочей среде — неправильный, поскольку результат зависит от непредсказуемого распределения времени между потоками.

Инструмент racecheck из compute-sanitizer анализирует обращения к общей памяти и сообщает об опасных ситуациях, которые обычный запуск скрывает.

Запись в общий массив с последующим чтением соседней ячейки без барьера создаёт гонку, поскольку соседняя ячейка может быть ещё не записана.

Вставьте __syncthreads между записью и чтением, чтобы каждый поток в блоке завершил запись до того, как любой поток начнёт чтение.

racecheck также обнаруживает опасность записи после записи, когда два потока записывают данные в одну общую ячейку без заданного порядка.

Самая распространённая опасность — чтение после записи: поток читает значение, которое другой поток ещё записывает, и получает устаревшие данные.

racecheck указывает тип опасности, а также два участвующих потока и общий адрес, чтобы вы точно знали, какого барьера не хватает.

racecheck сосредоточен на памяти __shared__ внутри блока; для глобальных гонок обычно требуется продуманная логика или атомарные операции.

Как и в случае с memcheck, компиляция с флагом -lineinfo позволяет racecheck указать точную строку исходного кода для каждой опасности.

После добавления барьера снова запустите racecheck. Отчёт о нулевом числе опасностей означает, что обращения к общей памяти теперь правильно упорядочены.

Вы использовали racecheck, чтобы найти опасности при работе с общей памятью, прочитали его отчёты и устранили проблемы с помощью барьеров __syncthreads. Ошибки синхронизации обнаружены. 🎯