Leçon 15 : Diagnostiquer et corriger un accès avec stride
Dans la leçon précédente, on a vu ce qu'est le coalescing : quand des threads consécutifs d'un warp touchent des adresses consécutives, une transaction les sert tous, et quand ils sautent avec un grand stride, le warp consomme de nombreuses transactions gaspillées. Maintenant, on apprend la compéten
Imagine 4 ouvriers qui cueillent des pommes dans des rangées d'arbres. Au début, chaque ouvrier court vers une colonne différente à un bout éloigné du verger — 4 longues courses. Si à la place chaque ouvrier prend la pomme suivante dans une seule rangée serrée, tous travaillent côte à côte en une seule passe. On n'a pas changé le verger — seulement qui-cueille-quoi.
- accès avec stride
- Un motif où les threads consécutifs sautent avec un grand stride (par exemple data[threadIdx.x * N]), donc leurs adresses sont très espacées et chacune tombe dans une transaction mémoire distincte.
- coalesced contre non coalesced
- Coalesced : les adresses des threads voisins sont contiguës, une transaction pour tout le warp. Non coalesced : les adresses sont dispersées, le warp consomme de nombreuses transactions, la plupart gaspillées.
- correction par remappage de l'indice
- La correction pour l'accès dispersé : faire en sorte que les threads voisins touchent des adresses voisines — donner au thread t l'élément base + t au lieu de base + t*N, tout en mémoire globale.
- ordre row-major
- Une disposition où les éléments de la même ligne sont voisins en mémoire. Parcourir par ligne fait que les threads voisins touchent des adresses voisines (coalesced) ; parcourir par colonne saute selon width (dispersé).
