Eliminer les Deep Copies

On l’a vu, le premier bénéfice de la mémoire unifiée est qu’elle simplifie le modèle mémoire du calcul hétérogène en supprimant la nécessité de réaliser des deep copies lorsque des structures de données sont utilisées dans les noyaux. Pour mémoire, c’est le transfert sur le GPU des structures de données contenant des pointeurs qui nécessite de réaliser ces fameuses deep copies. Prenons l’exemple de structure de données du listing 2. Pour utiliser cette structure sur le GPU, il faut d’abord la copier avec tous ses membres dans la mémoire du GPU, puis copier l’ensemble les données vers lesquelles elle pointe, puis mettre à jour les pointeurs. C’est exactement ce que réalise le code du listing 3, assez complexe assurément, simplement pour passer un élément de structure à un noyau.

Comme vous pouvez l’imaginer, la lourdeur de ce code, qui ne fait que partager des structures de données complexes entre CPU et GPU, a un impact significatif sur la productivité. A contrario, le simple fait d’allouer notre structure dataElem dans la mémoire unifiée rend inutile tout le code préparatoire. Au final, seul demeure le lancement du noyau, qui utilise maintenant le même pointeur que celui du système hôte comme le montre le listing 4. Un progrès qu’on peut raisonnablement qualifier de significatif !

Partager des listes chaînées entre CPU et GPU

Ce progrès ne se limite pas à la simplification de la programmation : il autorise également des approches algorithmiques jusqu’ici inenvisageables en CUDA. Prenons un second exemple mettant en musique des listes chaînées.

Ces listes sont des structures très courantes mais, parce qu’elles contiennent des données imbriquées les unes dans les autres via des pointeurs, les transférer d’un espace mémoire à un autre s’avère généralement assez complexe. Au point que, sans la mémoire unifiée, le partage de telles listes est quasiment impossible.

La seule option consiste à allouer la liste dans la mémoire Zero-Copy (pinned host memory) mais, alors, les accès du GPU à la mémoire du CPU sont limités par la performance du bus PCI-Express. En allouant les données des listes chaînées dans la mémoire unifiée, le code GPU peut maintenant traverser ces listes en utilisant simplement leurs pointeurs, et ce avec toute la puissance du système mémoire du GPU. L’application ne gère plus qu’une seule liste, de sorte que des éléments peuvent y être ajoutés ou supprimés depuis le CPU ou depuis le GPU.

Porter un code existant contenant des structures de données complexes était jusque-ici un exercice réservé aux programmeurs d’un certain niveau. La mémoire unifiée le rend beaucoup plus simple pour tout le monde, ce qui devrait en pratique se traduire par des gains de productivité non négligeables.

La mémoire unifiée avec C++

C’est avec les structures C++ que la mémoire unifiée prend finalement tout son sens. Déjà, C++ simplifie le deep copy grâce à ses constructeurs de copie. Un constructeur de copie, pour mémoire, est une fonction qui sait comment instancier une classe, allouer l’espace de ses membres puis copier leur valeur depuis un autre objet. Par ailleurs, C++ autorise aussi la surcharge des opérateurs new et delete, si bien que, comme dans le code figurant au listing 5, nous pouvons créer une classe de base, appelée Managed, qui utilise cudaMallocManaged() dans l’opérateur new surchargé.

On peut dès lors faire en sorte que la classe String hérite de la classe Managed et implémente un constructeur de copie pour allouer et copier une chaîne de caractères dans la mémoire unifiée, comme au listing 6. De la même façon, la classe dataElem hérite de celle Managed, comme le montre le listing 7 : les classes C++ allouent leur espace de stockage dans la mémoire unifiée et les deep copies sont gérées automatiquement. Notez enfin qu’il est possible d’allouer un dataElem dans la mémoire unifiée de la même manière qu’avec tout autre objet C++ :

dataElem *data = new dataElem;

Attention, pour éviter les “trous” dans le mappage mémoire, vous devez vous assurer que chaque classe de l’arborescence hérite bien de la classe Managed. C’est même une règle d’or : tout ce qui est partagé entre CPU et GPU doit hériter de Managed. Vous pouvez aussi surcharger les opérateurs globaux new et delete afin que toutes les allocations soient réalisées dans la mémoire unifiée. Ce qui n’a évidemment de sens que si toutes les données sont utilisées par le GPU, faute de quoi elles seront transférées inutilement.

Dans le même ordre d’idée, il est désormais possible de passer un objet à un noyau soit par valeur soit par référence, comme illustré au listing 8. Que l’on travaille sur des deep copies ou des passages par valeur ou par référence, avec la mémoire unifiée tout fonctionne de la même façon. C’est ce qui rend les codes C++ appropriés à une exécution sur GPU.

Les indicateurs au vert

Ces quelques exemples synthétiques ne sont qu’un avant-goût du potentiel quasiment illimité qu’offre la mémoire unifiée. Cette première implémentation a pour vocation de rendre la programmation CUDA plus simple pour les scientifiques on informaticiens. Grâce à cudaMallocManaged(), un pointeur unique permet de partager des structures de données C et C++ complexes entre CPU et GPU. On peut donc désormais se concentrer sur l’écriture de noyaux applicatifs sans se préoccuper de la gestion des données et de la cohérence des copies entre le CPU et le GPU. Bien sûr, chacun reste libre d’utiliser cudaMemcpy() (et plus particulièrement cudaMemcpyAsync()) pour des questions de performances. Mais plutôt qu’une obligation, c’est maintenant devenu une histoire d’optimisation.

Les prochaines versions de CUDA augmenteront très certainement les performances de la mémoire unifiée en ajoutant des fonctions de préchargement de données et la possibilité d’annoter les transferts. De plus, l’architecture des GPU de la prochaine génération apportera un certain nombre d’améliorations en termes de performance et de flexibilité. Ainsi, l’allocateur système sera unifié, ce qui signifie que n’importe quelle mémoire (allouée avec malloc() ou sur la pile) pourra être partagée entre CPU et GPU. Par rapport aux limitations des premières versions de CUDA, quel chemin parcouru…