Qu'est-ce que Coalescence mémoire ?
Regroupement en un petit nombre d'accès mémoire des accès de threads voisins d'un warp à des adresses voisines. Aussi appelé accès coalescé.
Nom japonais: メモリコアレッシング
Domaine: Programmation GPU
Comment fonctionne la coalescence
La mémoire globale d'un GPU est accédée par morceaux de taille fixe, généralement des secteurs de 32 octets au sein de lignes de cache de 128 octets, et non mot par mot. Les threads s'exécutent par groupes appelés warps, généralement 32 threads émettant ensemble la même instruction de chargement. Lorsque le thread i d'un warp lit l'élément de 4 octets situé à base + 4i, les 32 requêtes tombent dans 128 octets contigus, et le matériel les combine en une ou quelques transactions mémoire. Cette combinaison est la coalescence mémoire, et elle permet au GPU d'utiliser presque tous les octets qu'il transfère.
Ce qui empêche la coalescence
Si les threads accèdent à des adresses avec un pas (stride), par exemple un élément sur quatre, le warp touche quatre fois plus de secteurs, et l'essentiel de chaque secteur chargé est jeté. Des accès aléatoires ou indirects peuvent exiger une transaction distincte pour chaque thread. Des adresses de départ mal alignées peuvent ajouter un segment supplémentaire. Dans chaque cas, la bande passante effective tombe à une fraction de la valeur crête, et comme de nombreux kernels GPU sont limités par la bande passante mémoire plutôt que par le calcul, des accès non coalescés ralentissent souvent un kernel d'un facteur important.
Tableau de structures ou structure de tableaux
La disposition des données décide en grande partie si les accès sont coalescés. Lorsque les positions de particules sont stockées sous forme de tableau de structures (AoS), les x, y et z de la particule 0 sont suivis de ceux de la particule 1 ; lorsque chaque thread lit x, les adresses sont espacées de 12 octets. Avec une structure de tableaux (SoA), toutes les valeurs x sont contiguës : des threads adjacents lisent des adresses adjacentes et les chargements du warp sont coalescés. La SoA est aussi généralement préférable pour la vectorisation SIMD sur CPU, pour la même raison : des voies voisines veulent des éléments voisins.
Techniques et confusion fréquente
Les corrections courantes consistent à associer l'indice de thread qui varie le plus vite, threadIdx.x en CUDA, à la dimension contiguë d'un tableau, à convertir des données AoS en SoA, et à faire transiter les données par la mémoire partagée sur puce, par exemple en chargeant une tuile de matrice par lectures de lignes coalescées puis en l'écrivant transposée. La coalescence concerne les transactions de la mémoire globale. Il ne faut pas la confondre avec les conflits de banc (bank conflicts) de la mémoire partagée, qui surviennent lorsque plusieurs threads d'un warp accèdent à des adresses différentes d'un même banc de mémoire partagée et que leurs accès sont sérialisés.
Questions fréquentes
Qu'est-ce que la coalescence mémoire en CUDA ?
C'est la combinaison, par le matériel, des accès à la mémoire globale des threads d'un warp en aussi peu de transactions que possible, ce qui se produit lorsque des threads adjacents accèdent à des adresses adjacentes et alignées.
Pourquoi une structure de tableaux est-elle préférable pour les accès mémoire sur GPU ?
Dans une structure de tableaux, chaque champ est stocké de façon contiguë : des threads adjacents lisant le même champ accèdent à des adresses consécutives et leurs chargements sont coalescés, alors qu'un tableau de structures introduit un pas entre eux.
Quelle est la différence entre coalescence mémoire et conflits de banc ?
La coalescence concerne la façon dont les accès à la mémoire globale d'un warp se traduisent en transactions mémoire, tandis que les conflits de banc concernent plusieurs threads accédant au même banc de la mémoire partagée sur puce, ce qui impose une sérialisation.
Termes associés
Bloc de threads Boucle grid-stride Conflit de banc CUDA Divergence de warp Masquage de latence Mémoire constante Mémoire épinglée (pinned memory) Mémoire partagée (shared memory) Mémoire unifiée