Pression sur les registres et déversements

Les registres sont le stockage le plus rapide dont dispose un thread, mais chaque SM n'en contient qu'un nombre limité. La quantité de registres utilisée par un noyau est appelée pression sur les registres.

Chaque thread résident puise dans un seul fichier de registres par SM. Plus chaque thread a besoin de registres, moins il est possible de maintenir de threads résidents simultanément.

Une utilisation élevée des registres limite directement le nombre de warps pouvant tenir sur un SM. Cette baisse d'occupation peut laisser au matériel trop peu de travail pour masquer la latence.

Lorsqu'un thread a besoin de plus de registres qu'il n'en existe, le compilateur déplace les valeurs supplémentaires. Ce dépassement est un déversement de registres vers une mémoire plus lente.

Les valeurs déversées vont dans la mémoire locale, située dans la DRAM hors puce. Chaque déversement remplace un accès rapide à un registre par un aller-retour lent.

Demandez au compilateur de signaler l'utilisation. L'ajout de --ptxas-options=-v à nvcc affiche le nombre de registres par thread ainsi que les octets déversés pour chaque noyau.

Le rapport répertorie les écritures et les lectures liées aux déversements. Toute valeur non nulle du nombre de déversements indique que votre noyau paie le coût du trafic lent vers la mémoire locale.

Vous pouvez définir une limite lors de la compilation. L'option --maxrregcount limite le nombre de registres par thread, au prix d'une légère perte de vitesse pour augmenter l'occupation.

Une indication propre au noyau est souvent préférable. __launch_bounds__ indique au compilateur la taille de votre bloc afin qu'il puisse budgéter les registres pour l'occupation souhaitée.

Souvent, il suffit de conserver moins de valeurs à la fois. Recalculer un résultat peu coûteux ou réduire la portée d'une variable diminue le nombre de registres qui restent actifs.

L'ILP et le déroulage des boucles augmentent la pression, tandis que les limites augmentent l'occupation. L'objectif est de trouver l'équilibre permettant l'exécution la plus rapide ; seul le profileur peut vous indiquer où il se trouve.

Vous avez appris qu'une forte pression sur les registres réduit l'occupation et peut provoquer des déversements vers une DRAM lente. Mesurez l'utilisation, limitez les registres et laissez le profileur vous guider. 🎯