Leçon 25 : Scan parallèle — le kernel de doublement
Dans la leçon précédente, on a vu ce que calcule une somme de préfixes et l'idée du doublement : à chaque tour, chaque élément s'ajoute le voisin situé à une distance d, et d double (1, 2, 4...). Ici, on le construit comme un vrai kernel CUDA. On travaille au sein d'un seul block avec un tableau en
Chaque thread possède une position dans le tableau partagé. À chaque étape, il lit d'abord la valeur située d places à sa gauche et la met de côté, attend que tous aient lu (__syncthreads), et alors seulement se l'ajoute, puis attend à nouveau que tous aient écrit avant l'étape suivante. Les attentes sont ce qui maintient l'ordre et évite l'écrasement.
- Hillis-Steele
- Le kernel de scan par doublement en log2(n) étapes : à l'étape d, chaque thread avec tid>=d ajoute la valeur de tid-d. d double à chaque étape.
- tile en mémoire partagée
- Un tableau __shared__ pour tout le block. Le scan y charge l'entrée, travaille sur place à travers toutes les étapes, et écrit finalement dans out.
- barrières lecture/écriture
- __syncthreads() après la lecture (pour que tous lisent avant que quiconque n'écrive) et après l'écriture (pour que les écritures soient visibles à l'étape suivante).
- nombre d'étapes log2(n)
- La boucle for(d=1; d<n; d*=2) s'exécute log2(n) fois, parce que d double à chaque étape jusqu'à atteindre n.
