Limites des registres et de la mémoire partagée

Chaque SM possède une réserve fixe de registres et de mémoire partagée. Chaque bloc résident doit prélever sa part dans ces réserves.

Votre noyau utilise un certain nombre de registres par thread. Multipliez ce nombre par celui des threads par bloc pour obtenir le coût en registres d’un bloc.

Si chaque thread nécessite beaucoup de registres, moins de threads peuvent tenir sur le SM, et moins de blocs restent résidents. Les noyaux gourmands en registres perdent en occupation.

Compilez avec cet indicateur pour que nvcc affiche le nombre de registres par thread et vous permette de repérer les noyaux trop gourmands en registres.

Vous pouvez imposer une limite avec une contrainte de lancement afin que le compilateur utilise moins de registres et permette à davantage de warps de rester résidents.

Chaque bloc peut demander de la mémoire partagée. Le SM ne peut accueillir que le nombre de blocs autorisé par sa réserve partagée, souvent comprise entre 48 et 100 KB.

Demandez une grande tuile __shared__ et seuls un ou deux blocs pourront tenir sur chaque SM. Les grandes tuiles échangent de l’occupation contre une meilleure réutilisation des données.

Le SM calcule le nombre de blocs autorisés par les registres, la mémoire partagée et la limite de warps. C’est la plus petite de ces valeurs qui détermine l’occupation.

Lorsqu’un thread a besoin de plus de registres que la limite autorisée, les registres supplémentaires débordent vers une mémoire locale lente. Ces débordements peuvent nuire davantage que la faible occupation.

Réduire le nombre de registres ou la mémoire partagée augmente l’occupation, mais une réduction trop agressive provoque des débordements. Le meilleur compromis repose sur l’équilibre.

Lisez toujours l’utilisation réelle des registres et de la mémoire partagée fournie par le compilateur avant d’optimiser. Deviner conduit généralement à modifier le mauvais paramètre.

Vous avez vu que les registres et la mémoire partagée constituent des budgets fixes du SM, et que la limite la plus stricte plafonne l’occupation. Surveillez les débordements. 🧮