Cercare le race con racecheck

Si verifica una data race quando due thread accedono alla stessa memoria condivisa e almeno uno scrive, senza alcun ordinamento tra le operazioni. 😬

Le race condition spesso producono il risultato corretto nei test e quello errato in produzione, perché il risultato dipende dalla temporizzazione imprevedibile dei thread.

Lo strumento racecheck di compute-sanitizer analizza gli accessi alla memoria condivisa e segnala i rischi che un'esecuzione normale nasconderebbe.

Scrivere in un array condiviso e poi leggere la posizione adiacente senza una barriera crea una race condition, perché la posizione vicina potrebbe non essere ancora stata scritta.

Inserisca __syncthreads tra la scrittura e la lettura, così ogni thread del blocco termina di scrivere prima che qualcuno legga.

racecheck rileva anche un rischio di write-after-write, in cui due thread scrivono nella stessa posizione condivisa senza un ordine tra loro.

Il rischio più comune è read-after-write: un thread legge un valore che un altro thread sta ancora scrivendo, quindi legge dati obsoleti.

racecheck indica il tipo di rischio, i due thread e l'indirizzo condiviso coinvolti, così può capire esattamente quale barriera manca.

racecheck si concentra sulla memoria __shared__ all'interno di un blocco; per le race condition globali servono generalmente una logica accurata o operazioni atomiche.

Come per memcheck, la compilazione con -lineinfo permette a racecheck di indicare la riga esatta del codice sorgente di ogni rischio.

Esegua nuovamente racecheck dopo aver aggiunto la barriera. Un report con zero rischi significa che gli accessi alla memoria condivisa sono ora ordinati correttamente.

Ha utilizzato racecheck per trovare i rischi nella memoria condivisa, letto i relativi report e risolto i problemi con barriere __syncthreads. I bug di temporizzazione sono stati portati alla luce. 🎯