C’est l’un des plus importants développements dans l’histoire de CUDA : les mémoires CPU et GPU sont dorénavant unifiées du point de vue logiciel. Comment adapter vos codes et quels en sont les bénéfices concrets pour la programmation d’applications parallèles ? C’est ce que nous allons voir en détail maintenant, exemples à l’appui.

Mark Harris, Ph.D. – Research Scientist, NVIDIA.

Que ce soit dans un PC ou un dans un nœud de cluster x86, les mémoires du CPU et du GPU, physiquement séparées, sont reliées par le bus PCI-Express. Jusqu’à la sortie de CUDA 6, c’est cette architecture que le programmeur se devait d’avoir en tête pour allouer les données à partager entre processeur principal et processeur(s) graphique(s), et donc pour gérer explicitement leurs transferts. Quiconque en a fait l’expérience le sait, cette structuration rendait la programmation CUDA plus complexe. En créant une zone de mémoire partagée, la mémoire unifiée permet maintenant de lier ce qui séparait jusque-là CPU et GPU. Avec un pointeur unique, cette zone mémoire est dorénavant accessible par l’un et par l’autre. Le point clé de ce système réside dans le transfert automatique des données allouées dans cette mémoire unifiée. Avec pour conséquence principale que le code s’exécutant sur le CPU la considère comme de la mémoire GPU, et vice-versa.

De ce fait, clairement, la mémoire unifiée simplifie beaucoup la gestion mémoire, comme l’illustre les listing 1a et 1b. Ces deux codes chargent un fichier puis trient les données qui sont ensuite utilisées par le CPU avant de libérer la mémoire. Le code de droite s’exécute sur le GPU en utilisant CUDA et la mémoire unifiée. La seule différence entre ces deux codes est que la version CUDA lance un noyau sur le GPU (avec synchronisation après exécution) et alloue de l’espace pour charger le fichier dans la mémoire unifiée en utilisant le nouvel appel cudaMallocManaged() de l’API.

Si vous avez déjà programmé en CUDA C/C++, la simplicité du code de droite ne vous échappera pas. Remarquez que la mémoire est allouée une seule fois et qu’elle est accessible par le CPU et le GPU au travers d’un pointeur unique. Les données du fichier sont d’abord chargées sur le CPU et leur pointeur est ensuite directement passé au noyau CUDA s’exécutant sur le GPU. Dès que le noyau a terminé son exécution, les données sont à nouveau accessibles par le CPU. En transférant automatiquement les données là où les accès ont lieu, le runtime CUDA masque toute la complexité des mouvements et transferts de données.

Un double bénéfice pour le développeur

Par rapport aux anciennes versions de CUDA, la mémoire unifiée apporte deux principaux avantages. Le premier, c’est une programmation et un modèle mémoire simplifiés. En faisant en sorte que la gestion mémoire soit une optimisation plutôt qu’une nécessité, la mémoire unifiée facilite la transition de la programmation classique à la programmation parallèle. Les développeurs peuvent maintenant se concentrer uniquement sur les noyaux “métier” sans avoir à se préoccuper des détails d’allocation des données et de gestion des copies. L’apprentissage de la programmation CUDA et le portage des codes existants vers les GPU y gagnent ainsi en facilité et en rapidité. Mais que les programmeurs chevronnés se rassurent, cette amélioration ne concerne pas uniquement les débutants : les exemples que nous donnons ensuite montrent aussi comment la mémoire unifiée facilite l’utilisation de structures de données complexes, et comment elle livre toute sa puissance lorsqu’elle est utilisée en C++.

Le second avantage principal de la mémoire unifiée concerne la réconciliation toujours délicate entre performance et localité des données. En transférant à la demande les données entre le CPU et le GPU, la mémoire unifiée offre le même niveau de performance que si les données résidaient sur le GPU, tout en les rendant partageables globalement. La complexité de ce mécanisme est masquée par le pilote et par le runtime CUDA dont l’objectif commun est d’utiliser au maximum la bande passante de chacun des processeurs. Les 250 GB/s de la GDDR5 sont en effet vitaux pour alimenter la puissance de calcul d’un GPU Kepler, c’est-à-dire pour éviter qu’il n’attende sans rien faire que les données lui parviennent.

Il est important de noter qu’un programme CUDA utilisant les streams et cudaMemcpyAsync() pour recouvrir efficacement les calculs et les transferts de données restera plus performant qu’un code utilisant la mémoire unifiée. A cela, rien que de très normal : le runtime CUDA n’a pas connaissance des intentions du développeur quant à l’endroit et au moment où les données seront utilisées.

Pour en optimiser la gestion et la concurrence CPU-GPU, un développeur CUDA peut toujours réaliser des allocations explicites dans la mémoire du GPU et des transferts asynchrones. La mémoire unifiée doit à ce titre être conçue avant tout comme une fonction de productivité accrue pour quiconque désire se lancer dans la programmation accélérée. Autrement dit, son implémentation ne supprime pas les autres fonctions avancées de CUDA.

Mémoire unifiée ou adressage virtuel unifié ?

Depuis la version 4, CUDA supporte l’adressage virtuel unifié (UVA). Mais, point important, si la mémoire unifiée repose sur l’UVA, ces deux mécanismes ne sont pas identiques. L’UVA offre un espace d’adressage unique sur toute la mémoire du système et permet ainsi à un code GPU d’utiliser des pointeurs sur des données situées dans une zone mémoire quelconque, du GPU ou du CPU, ou même dans la mémoire partagée du système.

L’UVA permet aussi d’utiliser cudaMemcpy() sans avoir à spécifier où sont les paramètres d’entrée et de sortie. Enfin, l’UVA autorise les “Zero-Copy” pour des données allouées dans la mémoire pinned du CPU, elle-même directement accessible par le GPU au travers du bus PCI-Express sans devoir réaliser un memcpy. Les Zero-Copy sont tout aussi pratiques que la mémoire unifiée mais n’offrent pas le même niveau de performances, à cause de la faible bande passante du bus PCI-Express et de sa latence élevée.

Au contraire de la mémoire unifiée, l’UVA ne réalise pas automatiquement les transferts de données d’un emplacement à un autre. La mémoire unifiée, en revanche, est capable de transférer tout une page mémoire ce qui, pour l’anecdote, a nécessité un important travail d’ingénierie au niveau du runtime et du driver CUDA. Les exemples que nous allons voir maintenant illustrent ce que cette avancée rend possible.