Memburu Perlombaan dengan racecheck

Perlombaan data terjadi ketika dua thread mengakses memori bersama yang sama dan setidaknya salah satunya menulis, tanpa pengurutan di antara keduanya. 😬

Perlombaan sering menghasilkan jawaban yang benar saat pengujian dan jawaban yang salah dalam produksi karena hasilnya bergantung pada waktu thread yang tidak dapat diprediksi.

Alat racecheck dari compute-sanitizer menganalisis akses memori bersama dan melaporkan bahaya yang mungkin tersembunyi dalam proses normal.

Menulis ke larik bersama lalu membaca slot milik tetangga tanpa penghalang menimbulkan perlombaan, karena slot tetangga mungkin belum ditulis.

Sisipkan __syncthreads di antara penulisan dan pembacaan agar setiap thread dalam blok selesai menulis sebelum thread mana pun membaca.

racecheck juga menangkap bahaya tulis setelah tulis, ketika dua thread menyimpan data ke slot bersama yang sama tanpa pengurutan di antara keduanya.

Bahaya yang paling umum adalah baca setelah tulis: sebuah thread membaca nilai yang masih ditulis thread lain, sehingga membaca data lama.

racecheck menyebutkan jenis bahaya beserta kedua thread dan alamat bersama yang terlibat, sehingga Anda tahu penghalang mana yang hilang.

racecheck berfokus pada memori __shared__ di dalam blok; perlombaan global biasanya memerlukan logika yang cermat atau operasi atomik.

Sama seperti memcheck, kompilasi dengan -lineinfo memungkinkan racecheck menunjuk baris sumber yang tepat untuk setiap bahaya.

Jalankan ulang racecheck setelah menambahkan penghalang. Laporan nol bahaya berarti akses memori bersama Anda kini telah diurutkan dengan benar.

Anda menggunakan racecheck untuk menemukan bahaya pada memori bersama, membaca laporannya, dan memperbaikinya dengan penghalang __syncthreads. Bug waktu pun terungkap. 🎯