Leçon 17 : __syncthreads() et conditions de course

Quand tu charges des données en mémoire partagée puis que tu les lis entre différents threads, un problème apparaît : les threads d'un block ne tournent pas exactement au même rythme. Un thread peut arriver à la ligne de lecture avant qu'un autre thread ait fini d'écrire la case dont il a besoin — c

Un groupe de personnes range des livres sur une étagère partagée, et ensuite chacune a besoin de prendre un livre qu'une autre personne a rangé. Si tu commences à prendre des livres avant que tout le monde ait fini de ranger, certains tendront la main vers un espace vide. __syncthreads(), c'est dire « personne ne prend rien tant que tout le monde n'a pas rangé son livre ».

__syncthreads(): Une barrière au sein d'un block : tous les threads s'arrêtent à ce point jusqu'à ce que tous soient arrivés, et alors seulement ils continuent ensemble.
condition de course (race condition): Quand le résultat dépend de l'ordre dans lequel les threads arrivent à une écriture/lecture. Ici : lire la mémoire partagée avant que l'écriture soit terminée.
barrière (barrier): Un point d'attente que tous les threads du block doivent atteindre avant que l'un d'eux continue. Ça bloque les courses entre une écriture et une lecture.
portée de la barrière (block): __syncthreads() ne fonctionne qu'au sein d'un seul block. Elle n'a aucun effet entre différents blocks, et ne doit pas être appelée dans une branche divergente.