Cómo detectar condiciones de carrera con racecheck

Se produce una carrera de datos cuando dos hilos acceden a la misma memoria compartida y al menos uno escribe en ella, sin ningún orden entre sus operaciones. 😬

Las carreras suelen producir el resultado correcto en las pruebas y uno incorrecto en producción, porque el resultado depende de la temporización impredecible de los hilos.

La herramienta racecheck de compute-sanitizer analiza los accesos a la memoria compartida e informa de los conflictos que una ejecución normal ocultaría.

Escribir en un array compartido y después leer la posición de un vecino sin una barrera crea una carrera, ya que es posible que la posición del vecino aún no se haya escrito.

Inserte __syncthreads entre la escritura y la lectura para que todos los hilos del bloque terminen de escribir antes de que cualquiera lea.

racecheck también detecta un conflicto de escritura tras escritura, en el que dos hilos escriben en la misma posición compartida sin ningún orden entre ellos.

El conflicto más habitual es la lectura tras escritura: un hilo lee un valor que otro todavía está escribiendo, por lo que obtiene datos obsoletos.

racecheck indica el tipo de conflicto, los dos hilos implicados y la dirección compartida, para que sepa exactamente qué barrera falta.

racecheck se centra en la memoria __shared__ dentro de un bloque; las carreras en memoria global suelen requerir una lógica cuidadosa o atómicas.

Al igual que memcheck, compilar con -lineinfo permite que racecheck señale la línea exacta del código fuente de cada conflicto.

Vuelva a ejecutar racecheck después de añadir la barrera. Un informe con cero conflictos significa que los accesos a la memoria compartida están ahora correctamente ordenados.

Ha usado racecheck para encontrar conflictos en la memoria compartida, ha leído sus informes y los ha corregido con barreras __syncthreads. Los errores de temporización han quedado al descubierto. 🎯