Où l’on démontre que l’utilisation de mémoire globale est à la fois inutile et coûteuse en performances. A contrario, la mémoire partagée ne présente que des avantages.  Voici pourquoi, et surtout comment  l’utiliser efficacement…

Nous entamons ce premier épisode de l’atelier CUDA avec une technique d’optimisation fine des accès mémoire de vos applications – technique dont le bénéfice sera effectif sur la quasi-totalité des exécutables CUDA, quel que soit leur champ applicatif propre. Chimie, ingénierie, mécanique… nous sommes prêts à parier que son adoption se ressentira de façon mesurable sur les performances réelles de vos codes sources. Et cela, indépendamment de la génération d’accélérateur – Fermi ou Kepler – dont vous disposez.

Pour planter le décor, rappelons  qu’il existe plusieurs types de mémoire sur un GPU CUDA, chacune ayant une portée, une durée de vie et une stratégie de cache différente. Incontournable, la mémoire dite globale est traditionnellement utilisée pour les échanges avec le système et pour les entrées/sorties à partir des noyaux. Le qualificatif « global » se réfère ici à la portée, car l’espace exposé est consultable et modifiable à la fois par l’hôte et par l’accélérateur.

La mémoire globale peut être déclarée via une variable de portée globale, en utilisant le mot-clé __device__ comme à la première ligne du listing 1, ou allouée dynamiquement via une instruction cudaMalloc() puis affectée à un pointeur C standard comme à la dernière ligne du même listing. Dans les deux cas, elle persiste pendant toute la durée de vie de l’application et, selon les capacités propres de l’accélérateur, peut être cachée ou non.

Regroupements coupables

Jusque-là, rien d’extraordinaire, à ceci près que, dans une optique de performances, il faut bien être au clair avec le modèle d’exécution CUDA. Dans celui-ci, les threads sont d’abord regroupées en blocs affectés aux multiprocesseurs de l’accélérateur. Ensuite, pendant l’exécution, CUDA opère un regroupement plus fin des threads en “warps” de 32 threads maximum. Pourquoi ? Pour que le GPU exécute les instructions de chaque warp en mode SIMD et qu’intervienne effectivement la vectorisation, synonyme de parallélisation performante.

Si le regroupement de threads en warps est pertinent pour les opérations de calcul, il l’est au moins autant pour les accès mémoire. Grâce à lui, en effet, le GPU est capable de fusionner les ordres de chargement et de stockage de données issus des threads d’un même warp, de façon à diminuer au maximum le nombre de transactions, donc les besoins en bande passante. Notez à ce propos que pour les compute capabilities CUDA inférieures à 2.0 (carte Tesla C2050), la fusion s’opère sur des warps de 16 threads seulement.

Une expérience curieuse

Pour mieux nous rendre compte de ce qui se passe effectivement à l’usage et apprécier les bénéfices potentiels de la chose, faisons une petite expérience. Le listing 2 montre un petit noyau d’incrémentation auquel sont appliqués deux modèles. Le premier implémente un offset synonyme d’accès non-alignés au tableau de données. Le second, plus orthodoxe, implémente des accès alignés (“strided “). Le code peut être exécuté en simple ou en double précision selon que l’on passe le paramètre “fp64” en ligne de commande, chaque noyau prenant deux arguments : le tableau cible et un entier représentant l’offset d’accès aux éléments. Les noyaux sont appelés de façon itérative sur une large plage d’offsets.

A l’exécution, les tableaux alloués en mémoire globale sont alignés sur des segments de 256 octets par le pilote CUDA. Le GPU y accède par transactions de 32, 64 ou 128 octets, si bien que des accès mal alignés occasionnent une multiplicité de transactions. En cela, rien de très surprenant, à un point près. La baisse de bande passante n’est proportionnelle au taux de données pertinentes par transaction que sur les cartes Tesla les plus anciennes. A partir d’un compute capability de 2.0 (Tesla C2050, rappelons-le), l’ensemble des transactions gérées par le pilote couvre les besoins de la requête initiale, de sorte que la perte en performances est moins sensible. On paye encore clairement le défaut d’alignement, du fait notamment des transferts de données inutiles, mais dans une moindre mesure. Pourquoi ? Parce que sur les GPUs les plus récents, chaque multiprocesseur dispose d’un cache L1 de 128 octets. Grâce à lui, les accès mémoire initiés par les threads d’un même warp vont pouvoir être réunis dans un espace ultra-performant. Surprenante conclusion : avec les accélérateurs Tesla récents, l’effet de l’alignement reste finalement assez négligeable pour les accès mémoire séquentiels entre threads.

Pour les accès alignés, les résultats obtenus avec notre petit modèle de test sont conformes à nos attentes : plus la taille des blocs est élevée, plus la bande passante diminue. Techniquement, l’explication est simple. Quand des threads concurrentes accèdent de façon simultanée à des adresses éloignées en mémoire physique, les probabilités d’optimisation matérielle automatique sont quasi-nulles. Il faut simplement savoir que, plus l’accélérateur utilisé est récent, plus la dégradation des performances est linéaire, grâce à l’alignement des segments sur les capacités mémoire des cœurs.

Méchante mémoire globale

Que retirer de ces deux résultats ? La conclusion que l’utilisation de la mémoire globale n’est probablement pas la meilleure approche en termes de performances pures. D’où l’importance de perdre ce réflexe de pensée évident, et d’envisager la structure globale des accès mémoire autrement.

C’est là qu’intervient la mémoire dite « partagée ». Ce qualificatif, elle le doit au fait que toutes les threads d’un même bloc accèdent à l’espace de façon uniforme. C’est très pratique quand, par exemple, il est nécessaire que les threads indexent les plus hautes dimensions d’un tableau multidimensionnel – ce qui, par parenthèse, garantit l’alignement des accès. On aura donc intérêt à l’utiliser le plus souvent possible pour y stocker des blocs 2D constituant le tableau multidimensionnel avec lequel on travaille. Une fois cet espace alloué, les threads contigües peuvent y accéder de façon alignée sans perte de bande passante, contrairement à ce que nous venons d’observer pour la mémoire globale.

Des caches personnalisés

Parce que sa logique réside directement en hard, la mémoire partagée est bien plus rapide que les espaces logiques locaux ou globaux. Dès lors qu’il n’y a pas de conflit dans l’utilisation programmée des banques, son niveau de latence est environ 100 fois inférieur à celui d’un espace non caché. En pratique, la mémoire partagée est allouée par blocs de threads, si bien que toutes les threads d’un même bloc y ont accès directement. Ainsi, typiquement, des threads seront chargées de l’extraction de données depuis la mémoire globale, cependant que d’autres (du même bloc) seront dédiées aux calculs. Grâce à la synchronisation automatique des threads, il devient donc possible d’implémenter assez simplement des systèmes de caches dynamiques finement adaptés au type de données traitées. C’est d’ailleurs un des secrets de la programmation parallèle efficace.

La légère difficulté est ici d’éviter que ne se crée une situation de compétition (“race condition”). L’effet est bien connu, lorsque les threads d’un bloc s’exécutent logiquement en parallèle, rien ne garantit qu’elles s’exécutent physiquement en même temps, d’où la possibilité que les données partagées ne soient pas elles-mêmes synchronisées. La bonne nouvelle, c’est que cette synchronisation s’obtient sans douleur. CUDA offre pour cela une primitive simple, __syncthreads() : elle garantit qu’une thread ne peut poursuivre son exécution qu’à condition que toutes les autres threads du bloc soient elles-mêmes synchronisées. L’usage veut qu’on invoque cette primitive après que les données aient été stockées en mémoire partagée mais avant toute consommation. Bien observer cette rigueur logique a une conséquence positive : si toutes les threads appellent __syncthreads() au même moment de leur exécution, on évite toute possibilité de blocage du code à l’exécution.

Une technique très flexible

Le listing 3 montre un exemple concret de synchronisation, réutilisable dans quasiment tous les codes sources. La façon la plus simple de créer de la mémoire globale est d’utiliser le mot-clé __shared__ dans la déclaration de la variable mémoire. Comme le montre le listing, il existe toutefois d’autres approches parfois mieux adaptées au contexte, selon par exemple que l’on connaît l’espace disponible ou nécessaire au moment de la compilation ou au moment de l’exécution.

Techniquement, le code du listing 3 ne fait rien d’autre qu’inverser, en mémoire partagée, les données d’un tableau de 64 éléments. Mais l’approche n’est pas la même selon le mode de déclaration des tableaux et l’invocation des noyaux. L’utilisation de mémoire partagée statique est la plus évidente. Si l’on connaît la taille du tableau à la compilation (voir le noyau staticReverse()), il suffit de déclarer sa dimension le plus simplement du monde. Les indices original et inverse sont ensuite représentés par t et tr, respectivement, après quoi les threads peuvent copier les données depuis la mémoire globale via l’instruction s[t] = d[t], puis réaliser l’inversion deux lignes plus bas via d[t] = s[tr]. Entre les deux, la présence de la primitive __syncthreads() garantit qu’aucune opération de lecture de données partagées n’aura lieu avant que celles-ci aient été chargées dans leur état de consommation algorithmique.

Avantages induits

Cette approche offre un gros avantage avec les GPU d’ancienne génération. La fusion des opérations de lecture et d’écriture s’effectue de façon optimale parce que les accès en mémoire globale s’effectuent via l’index t, qui est lui-même linéaire et aligné. L’index inverse tr n’est utilisé que pour les accès en mémoire partagée, de sorte qu’on ne paye à l’exécution aucune restriction d’accès séquentiel.

Les trois autres noyaux du listing 3 utilisent de la mémoire partagée allouée dynamiquement – cas typique lorsque l’espace nécessaire n’est pas connu à la compilation. Dans de tels contextes, la taille d’allocation par bloc de threads doit être spécifiés (en octets) via un troisième paramètre de configuration, comme dans l’instruction :

dynamicReverse<<<1, n, n*sizeof(int)>>>(d_d, n);

Par ailleurs, remarquez que le noyau de mémoire dynamique, dynamicReverse(), déclare la mémoire partagée par le biais d’un tableau extern non dimensionné, selon la syntaxe :

extern __shared__ int s[];

(notez les crochets vides et l’utilisation du spécificateur extern).

Cette relative économie de moyens déclaratifs permet tout de même à CUDA de déterminer la taille désirée au lancement du noyau. Concernant le reste du code, il n’y a pas de différence avec le noyau staticReverse(), ce qui démontre la flexibilité de cette technique.
Pour finir sur ce point, un truc éprouvé. Si vous avez besoin de plusieurs tableaux dynamiques dans le même noyau, évitez de multiplier les entités. Il suffit de déclarer seul tableau extern non-dimensionné, puis d’utiliser des pointeurs pour le diviser en sous-tableaux, comme au listing 4-A. Ensuite, pour lancer le noyau, il vous faudra simplement spécifier la taille globale dont vous avez besoin, comme au listing 4-B. Bon développements !