Le pixel, le polygone et la matrice - le calcul par processeur graphique

1. GPGPU : le processeur graphique devient généraliste

1.1 Pourquoi calculer avec un GPU

Le GPU (Graphic Processing Unit - processeur graphique) est une puce spécialisée dans la gestion de l'affichage. A l'origine, il s'agissait uniquement de générer un signal vidéo analogique à partir du contenu d'une zone mémoire précise, plus ou moins réservée à cet usage. Au fur et à mesure du développement technologique, on a ajouté dans ces circuits spécialisés des unités de traitement dédiées aux opérations communément requises pour l'affichage graphique : tracé de droites, de cercles ou de rectangles. La prise en charge de la 3D, tout d'abord déléguée à des cartes additionnelles spécialisées (par exemple : 3DFX [1]), a naturellement été confiée par la suite au processeur graphique. Mais le traitement d'objet 3D, texturé ou non, nécessite des calculs mathématiques assez conséquents, donc les processeurs graphiques sont devenus de véritables coprocesseurs mathématiques, plutôt spécialisés dans le calcul vectoriel dans des espaces à trois voire quatre dimensions.

Au final, l'architecture d'un processeur graphique est devenue aussi raffinée que celle des processeurs généralistes habituels (Pentium, PowerPC ou autres) et le niveau d'intégration, le nombre de transistors, la finesse de gravure, la fréquence sont devenus comparables. Il reste toutefois des spécificités conçues pour optimiser les performances des opérations graphiques et de calcul 3D qui restent leur usage principal.

Un processeur généraliste comporte de nos jours de nombreuses unités de traitement plus ou moins spécialisées, plusieurs niveaux de mémoires caches globales, l'ensemble forme une structure assez complexe qui peut être dupliquée totalement ou partiellement dans des cœurs distincts, actuellement moins d'une dizaine. En comparaison, l'unité de calcul de base d'un processeur graphique est plus simple. En revanche, plusieurs dizaines de ces unités de calcul sont intégrées sur une seule puce graphique. Ce haut niveau de parallélisme créé initialement pour traiter simultanément la projection de plusieurs textures ou shaders est devenu disponible pour un tout autre usage avec l'apparition de kits de développement dédiés à cet effet : bienvenue dans l'ère du GPGPU (General-Purpose Processing on Graphic Processing Units).

1.2 NVidia CUDA

NVidia a naturellement suivi cette démarche et participé complètement à cette évolution technologique qui lui offre de nouvelles opportunités de marché et la possibilité de révéler un véritable savoir faire dans le domaine du calcul haute performance. Ce qui ajoutera inévitablement à sa crédibilité si jamais, d'aventure, l'ambition de concurrencer les fondeurs de processeurs généralistes devenaient son nouveau cheval de bataille. Une nouvelle ligne de produits a été introduite spécialement pour le calcul haute performance : la famille de puces Tesla. Elle s'accompagne d'un kit de développement dédié qui s'articule autour de CUDA (Compute Unified Device Architecture), qui permet de rédiger le traitement à confier au GPU dans un langage très proche du C. La mise en œuvre est facilitée par l'API fournie qui offre tous les services d'initialisation et d'invocation nécessaires. NVidia propose même un compilateur C modifié, baptisé nvcc, qui permet d'écrire ses premiers programmes CUDA le plus simplement possible. Il traite des fichiers sources qui contiennent aussi bien le pseudo-code C à destination du GPU que le code C à destination du processeur de l'hôte, qui reste nécessaire pour la prise en charge de toutes les entrées/sorties et interactions avec l'utilisateur. En effet, si le GPU peut maintenant réaliser toutes sortes de calcul, il n'a pas accès à toute la plate-forme hôte et ne peut finalement rien faire d'autre que du calcul (il peut néanmoins afficher des choses si le résultat de ces calculs est une image ou un modèle 3D). Il sera donc impossible d'ajouter le moindre printf dans le code GPU pour voir ce qu'il s'y passe. Ceci constitue d'ailleurs une difficulté spécifique pour le développeur car il est vain de vouloir inclure des traces dans le code GPU, ce qui ne permet pas un débogage facile. Les dernières versions du kit incluent un débogueur spécialisé pour contourner cette difficulté.

Le kit de développement (SDK) est librement téléchargeable depuis le site CUDA zone [3]. Sont publiées sur ce site les versions Windows, Mac OSX et Linux (distribution RedHat, Ubuntu et SUSE, 32 et 64 bits) du kit, la documentation et les drivers propriétaires pour les cartes graphique. Ces derniers sont obligatoires pour avoir accès aux fonctions CUDA, l'alternative libre n'est pas (encore ?) capable de prendre en charge les extensions CUDA. En pratique, le kit est fourni sous forme d'archive auto-extractible et devrait fonctionner sur toute distribution compatible avec celles officiellement prises en charge, c'est-à-dire construite sur les mêmes versions des bibliothèques système (libc principalement) et pour lesquelles le driver propriétaire est disponible.

L'exemple ci-dessus est relativement trivial dans la mesure où les vecteurs traités sont suffisamment petits pour être pris en charge en une seule fois. Dans le cas contraire, il faut préciser l'organisation de la grille de blocs (selon la terminologie CUDA), qui réaliseront les calculs par paquet à l'aide de l'opérateur <<<dim_grille,dim_blocs>>>.

CUDA a été intégralement développé par NVidia et c'est évidemment une technologie propriétaire. Le seul engagement de NVidia est d'inclure la compatibilité CUDA à toutes ses futures puces graphiques. Si CUDA, de par son potentiel, mérite toute notre attention, son caractère propriétaire n'offre pas la moindre pérennité à un développement libre. De plus, le langage et les API n'ayant fait l'objet d'aucune normalisation, la pérennité de tout logiciel écrit en CUDA, même propriétaire, dépend exclusivement de la bonne volonté de NVidia. Heureusement, la concurrence n'est pas restée les bras croisés et le khronos group a publié fin 2008 les spécifications de la norme OpenCL [4]. La présentation de CUDA se clôt ici. En effet, OpenCL et CUDA sont suffisamment similaires pour que les principes généraux de l'un soient applicables à l'autre sans trop d'efforts. Une bonne expérience d'OpenCL et la documentation CUDA [5] seront largement suffisantes pour travailler correctement là où des contraintes externes interdisent de travailler avec un autre environnement. A partir de la version 3 du SDK CUDA, le support d'OpenCL est inclus.

1.3 ATI Stream

Evidemment, ATI (ou devrais-je dire désormais, AMD) ne pouvait pas laisser NVidia seul sur ce marché et a rapidement proposé son offre alternative : ATI Stream [6]. Le choix a été fait ici de proposer directement une mise en œuvre d'OpenCL, non seulement pour les puces graphiques ATI, mais aussi pour les processeurs x86 dotés des jeux d'instructions SSE3. Ainsi, même en l'absence de carte graphique compatible, ce kit reste utilisable avec n'importe quel PC récent et sera à lui seul suffisant pour s'initier à OpenCL dans de bonnes conditions, c'est-à-dire sans devoir passer par l'utilisation d'un émulateur. Le kit est disponible pour Windows et Linux (RedHat, Ubuntu et SUSE, 32 et 64 bits) et comme son concurrent, il doit pouvoir s'installer sur toute distribution construite sur les mêmes versions de bibliothèques système. Concernant les puces graphiques, là encore, le driver adéquat est nécessaire. Il s'agit du driver Catalyst qui, lui aussi, est encore propriétaire malgré les bonnes intentions affichées par AMD.

Pour l'installation du SDK ATI Stream [7], AMD a fait un service minimum, puisqu’il ne fournit qu'une archive tar compressée et l'utilisateur devra lui-même adapter manuellement l'installation à sa distribution pour mettre à jour des variables d'environnement et créer des liens symboliques afin d'autoriser un chargement correct des bibliothèques dynamiques.

1.4 OpenCL et GPGPU

OpenCL [8] est issu de OpenGL. Pour qui connaît bien cette dernière API, le lien de parenté est visible dans la structure de l'API proposée et la définition d'un code OpenCL ressemble beaucoup à celle d'un shader OpenGL. Le principal intérêt d'OpenCL est qu'il tente d'abstraire au mieux les ressources de calcul quelles qu'elles soient. Le code rédigé doit donc être très portable, d'un système de calcul composite à un autre, qu'il soit à base de GPU ou d'autre chose (par exemple, le processeur Cell BE qui équipe la console Playstation 3 [9]), à l'exception sans doute du code d'initialisation des périphériques et de leurs drivers spécifiques.

1.5 Compléments d'information

CUDA et OpenCL ne sont pas les seuls outils disponibles pour le GPGPU. Microsoft propose aussi sa solution, mais comme la société a maintenant de longue date tourné le dos à la norme OpenGL, il lui était difficile d'adopter OpenCL. Elle présente donc DirectCompute avec sa dernière version de DirectX.

Enfin, pour être complet, il faut noter que NVidia fait beaucoup d'efforts pour offrir des outils les plus proches possibles des besoins des utilisateurs traditionnels de gros moyens de calcul et les aider à migrer leurs applications vers CUDA avec un minimum de travail : CUDA Fortran (équivalent de CUDA C, mais basé sur le langage FORTRAN au lieu du C), bibliothèques CUBLAS (BLAS accéléré par GPU) et CUFFT (FFT accéléré par GPU), sont des extensions librement téléchargeables, comme le SDK lui-même.

2. Premiers pas en OpenCL

Avant toute chose réellement utile, quelques initialisations sont évidemment nécessaires. Un ensemble de fonctions permet d'interroger le système pour faire l'inventaire des périphériques de calcul, connaître leur caractéristiques et les configurer avant utilisation. Les ressources de calcul OpenCL sont réparties en plates-formes, chaque plate-forme pouvant être composée de plusieurs périphériques (device). Les fonctions clGetPlatformIDs et clGetPlatformInfo permettent respectivement de faire l'inventaire des plates-formes et de connaître leurs caractéristiques, tandis que les fonctions clGetDeviceIDs et clGetDeviceInfo font de même pour les périphériques. Une fois le périphérique et la plate-forme choisis, il faut, pour accéder à leurs services, créer un contexte à l'aide de clCreateContext. Lors de cette création, les paramètres définissant le mode de fonctionnement et d'utilisation souhaité sont définis. Dans l'exemple qui suit, on crée un contexte avec la configuration par défaut sur le premier périphérique OpenCL trouvé, sans aucune considération pour ses capacités et caractéristiques :

De façon analogue au shader OpenGL, le programme OpenCL est transmis au périphérique de calcul et compilé à l'aide d'un tampon de chaîne de caractères via des fonctions de l'API OpenCL :

- clCreateProgramWithSource : transmettre à l'API le buffer contenant le code source.

- clBuildProgram: compiler le programme. Dès que la compilation est terminée, la fonctionclGetProgramBuildInfopermet d'en connaître le compte-rendu, et en particulier d'obtenir les messages d'erreur en cas d'échec.

Pour terminer, il est nécessaire de créer une queue de commande et un noyau qui correspond à la fonction du code source OpenCL que l'on souhaite utiliser.

Avant de lancer le calcul, il faut d'abord allouer de la mémoire dans le périphérique de calcul et y copier les données à traiter, sans oublier d'allouer également la mémoire qui recevra les résultats.

Nous sommes alors enfin en mesure de lancer le calcul : dans un premier temps, on associe les tampons mémoire du périphérique à chacun des paramètres, puis on soumet l'appel de fonction à la queue de commande du périphérique.

Pour terminer le cycle, on demande la copie des résultats dans la mémoire de l'hôte et si les tampons mémoire du périphérique n'ont plus d'usage, il peuvent être libérés.

La compilation ne pose pas de problème particulier à partir du moment où les bibliothèques de développement OpenCL sont correctement installées : en-tête à inclure dans /usr/include/CL et bibliothèques dynamiques dans /usr/lib, à adapter selon les conventions de votre distribution préférée. Il suffit alors d'inclure libOpenCL.so à l'édition des liens.

3. OpenCL, la portabilité entre GPU et autres processeurs

ATI Stream n'étant que la mise en œuvre OpenCL par AMD, le SDK V3 de CUDA intégrant désormais la prise en charge d'OpenCL et disposant des deux systèmes compatibles chacun avec l'un des deux concurrents, j'ai pu vérifier la portabilité d'un même programme entre la plate-forme NVidia et la plate-forme ATI :

- Le premier système est un Zotac MAG HD ND-01, c'est-à-dire un nettop construit sur le couple Intel Atom 330/Nvidia ION. Il est animé par une distribution Ubuntu 09.04 avec le SDK CUDA V3 bêta (la version finale a été publiée pendant la rédaction de cet article).

- Le deuxième est un PC de bureau standard comprenant un processeur Intel Core 2 Duo et une carte graphique Radeon HD 4890. Ce système est aussi doté d'une distribution Ubuntu 09.04 accompagnée du SDK ATI Stream.

Sans tenir compte du temps d'exécution, car le GPU ION ne prétend pas rivaliser avec la puce ATI, la portabilité au niveau source est vérifiée pour le petit programme d'évaluation exposé ici. Il s'exécute à l'identique et les deux vecteurs sont correctement additionnés. Si on dispose sur chaque système de distributions compatibles au niveau binaire, le programme le sera aussi.

Pour observer un peu plus finement le comportement de chaque plate-forme, quelques essais complémentaires ont été menés. Tout d'abord, à l'aide des fonctions clGetPlatformInfo et clGetDeviceInfo, voici l'inventaire et quelques caractéristiques des moyens de calculs OpenCL trouvés sur chaque plate-forme :

Sur chaque système, on observe d'abord l'identification de la plate-forme OpenCL elle-même. Sur le Zotac, avec le SDK NVidia, la plate-forme ne comporte qu'un seul périphérique de calcul : le GPU. Par contre, sur la plate-forme ATI, nous voyons apparaître deux périphériques pris en charge : le GPU bien sûr, mais également le processeur principal, ce qui confirme les spécifications établies par AMD.

Nous pouvons ensuite observer ce qui est produit par la phase de construction du noyau OpenCL. En effet, une fois le programme construit, la fonction clGetProgramInfo, à l'aide de l'option CL_PROGRAM_BINARIES, permet de l'interroger et d'accéder à un « binaire ». Ce binaire dépend évidemment du matériel cible. Mais sur les deux plates-formes, ce que nous obtenons n'est pas exactement un binaire au sens habituel :

- Côté NVidia, j'ai eu la surprise de trouver un code assembleur LLVM [10]. Chacun appréciera l'attitude de cette société qui ne rechigne pas à utiliser le libre pour son propre usage et qui refuse toute ouverture concernant les données techniques de ses produits.

- Côté ATI, on obtient d'une part une simple indirection vers un vrai binaire au format ELF quand on cible le processeur Core 2, et un paquet binaire totalement opaque (vis-à-vis d'un examen trivial du bloc binaire reçu) quand on cible le GPU.

Ces comportements assez différents ne permettent pas d'envisager de réaliser un « pré-compilateur » OpenCL sur la simple base de la fonction clGetProgramInfo qui nous a permis d'obtenir ces binaires. En particulier, le cas de la cible Core 2 duo, de par son mode d'indirection vers un fichier temporaire sur /tmp, c'est-à-dire un espace potentiellement volatil, en limite la persistance. Ce qui est restitué ici par clGetProgramInfo n'est pas suffisamment fiable pour n'importe quel usage ultérieur. Ce sera quand même suffisant pour limiter la charge de la compilation au strict nécessaire. Après tout, ceci n'est réellement gênant que pour celui qui souhaite absolument cacher tous ses sources, ce qui est loin d'être le cas de l'audience cible du présent article.

4. Programmation parallèle

4.1 work-group et work-item

Pour faire simple, tout système de calcul OpenCL se divise hiérarchiquement en deux niveaux : work-group et work-item. A chaque niveau hiérarchique correspond un niveau de mémoire :

- Espace global : commun à tous les work-groups. Il s'agit des variables et paramètres d'un noyau déclaré avec l'attribut __global. Quand employé en tant que paramètre de noyau, cet espace permet de communiquer avec l'hôte (c'est-à-dire le partage de données entre l'hôte et le noyau). Chaque paramètre global doit alors être associé à un espace mémoire de l'hôte (voir l'utilisation de clCreateBuffer et clSetKernelArg dans les exemples ci-avant).

- Espace local : commun à tous les work-items d'un work-group. Il s'agit des variables et des paramètres d'un noyau déclarés avec l'attribut __local. Ces espaces sont inaccessibles pour l'hôte. Quand un noyau possède au moins un paramètre __local, l'hôte doit définir sa taille. C'est d'ailleurs le seul moyen d'allouer dynamiquement une variable locale. Pour ce faire, il faut utiliser clSetKernelArg, avec NULL pour quatrième argument, qui est normalement prévu pour indiquer l'espace mémoire de l'hôte qui est associé au paramètre.

- Espace privé : local à un work-item. Il s'agit des variables locales déclarées dans le corps d'un noyau.

Physiquement, ces espaces mémoire peuvent être répartis au plus près du besoin. Dans le cas d'un GPU, l'espace global sera la RAM de la carte graphique, synchronisée avec la mémoire du système hôte. L'espace local pourra être pris en charge par la mémoire cache de la puce, tandis que l'espace privé pourra être réparti directement sur les registres d'une unité de calcul (ou le cache de niveau inférieur s'il existe). Les performances réelles lors des accès à ces trois espaces dépendent des capacités propres de chaque équipement et de son driver OpenCL. Si l'espace global est mis en cache de façon transparente, les accès seront la plupart du temps aussi rapides que pour l'espace privé. Dans le doute, il vaut mieux bien choisir l'affectation des variables temporaires, surtout que l'affectation réelle est opaque car à la charge de chaque driver OpenCL. On réservera donc autant que possible l'usage de l'espace global à l'échange des données d'entrée et de sortie avec le système hôte.

Par principe, le découpage des travaux en work-group et work-item a pour but d'autoriser leur réalisation en concurrence, c'est-à-dire en parallèle. Le découpage et la répartition sont quasiment automatiques quand il est possible d'effectuer directement une parallélisation sur les données, c'est-à-dire quand un traitement rigoureusement identique (le même algorithme qui peut quand même inclure des conditions sur les données d'entrées) est appliqué à plusieurs données. Par contre, dès qu'une variable de sortie dépend d'une autre variable de sortie, il n'est plus possible d'atteindre le même niveau de parallélisme. Dans ce cas, il sera quand même possible d'accomplir une parallélisation sur les tâches pour réaliser en concurrence des algorithmes différents qui n'ont pas de liens entre eux.

- Parallélisation sur les données automatique : un même traitement appliqué à un grand nombre de données va être réparti de façon naturelle. Chaque élément de donnée est pris en charge par un work-item. Il suffit d'écrire un noyau qui prend en charge un unique élément et lors de son invocation, le nombre d'élément à traiter est précisé, ou plus exactement la taille de chaque dimension de l'espace des éléments. OpenCL se charge alors de créer les work-items, un pour chaque élément, et de les affecter simultanément aux unités de calcul disponibles. Dans le noyau, la fonction get_global_id permet d'obtenir le ou les indices de l'élément. En effet, il est possible de définir d’un à au moins trois indexes d'itération, pour extraire les éléments dans des structures d’une à au moins trois dimensions. Le nombre exact de dimensions possible peut être obtenu à l'aide de la fonction clGetDeviceInfo avec le paramètre CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS. La taille de chaque dimension sera précisée lors de l'invocation du noyau depuis le processus de l'hôte.

Ce mode de répartition est le plus simple à mettre en œuvre, mais la gestion déléguée des work-groups et work-items interdit toutes connaissance a priori de leur répartition et dans ce cas, le partage de l'espace local n'est pas établi a priori. Il n'est donc pas recommandé dans ce cas de travailler avec des variables __local et on se limitera par conséquent à l'espace global et à l'espace privé.

- Parallélisation sur les données contrôlé : le programmeur définit lui-même du découpage en work-group et work-item en indiquant précisément le nombre de work-items que contient chaque work-group dans chaque dimension. Le but est d'optimiser ce découpage selon les caractéristiques du matériel disponible et du problème à traiter. On garde ici toute la maîtrise sur la répartition du travail, mais cela demande un peu plus d'effort pour le programmeur. La taille d'un work-group doit être définie dans chaque dimension : le sixième paramètre de clEnqueueNDRangeKernel est un tableau d'entiers dont la longueur correspond au nombre de dimensions spécifiées par le troisième paramètre. Il faudra prendre garde à ne pas outrepasser les capacités de la plate-forme cible. Pour connaître les limites, la fonction clGetDeviceInfo est à nouveau le point d'accès incontournable. Chaque paramètre disponible est identifié par une constante. Ceux correspondant aux tailles limites des work-groups sont les suivants :

- CL_DEVICE_MAX_WORK_GROUP_SIZE : nombre maximal de work-items dans un work-group ;

- CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : nombre maximal de dimensions. La norme OpenCL précise que ce nombre est au moins trois, on pourra donc éviter d'interroger ce paramètre tant qu'on ne dépasse pas trois dimensions ;

- CL_DEVICE_MAX_WORK_ITEM_SIZES : taille maximum d'un work-group pour chaque dimension.

- Parallélisation sur les tâches : dans ce modèle, les indexeurs ne définissent plus une donnée à traiter, mais définissent plutôt une action à mener. En fait, la sémantique des indexes qui localisent chaque work-item dans chaque dimension est libre. On peut l'utiliser pour mettre en œuvre un parallélisme orienté sur des threads identifiés, pour lequel on considère que chaque work-item est un pseudo thread et l'indexation sera alors utilisée pour identifier chaque instance de thread. Il est également possible d'exécuter des noyaux différents en concurrence. La fonction clEnqueueNDRangeKernel n'est pas bloquante. Comme son nom le suggère, elle ne fait qu'envoyer une requête dans une file d'attente et l'hôte peut poursuivre son travail sans attendre que le noyau invoqué ait terminé.

Bien sûr, des approches hybrides sont possibles : puisqu'on dispose d'un système d'indexation à au moins trois dimensions, une dimension pourra être réservée à l'identification de l'action que le work-item doit réaliser, tandis que les deux autres localisent la donnée sur laquelle l'action est appliquée.

4.2 Types vectoriels

Des types vectoriels sont disponibles pour la plupart des types scalaires courants : char, int, float et même double. La taille d'un tel vecteur est 2, 4, 8, ou 16. Le nom du type vectoriel est simplement le nom du type de base concaténé à la taille, ainsi le type d'un vecteur d'entiers de dimension 4 est int4. Le type image correspondant en C pour l'hôte est préfixé cl_, c'est-à-dire cl_int4 pour l'exemple précédent.

Evidemment, comme toujours avec OpenCL, le gain de performance dépend de l'architecture des moyens de calcul physiques disponibles. Dans tous les cas, mieux vaut utiliser les types vectoriels autant que possible, puisque OpenCL prend lui-même en charge les ajustements nécessaires en fonction de la taille des registres et du jeu d'instruction disponible. Avec la plupart des GPU, optimisés pour le traitement 3D, les opérateurs 4D sont effectivement disponibles et une addition de deux float4 pourra être faite en une seule opération. Dans le même ordre d'idée, sur un processeur CellBE dont les unités SPU ont des registres vectoriels de 128 bits, on pourra traiter en une seule opération une addition sur deux vecteurs de 16 octets, alors qu'il en faudra au moins deux pour additionner deux vecteurs double de dimension 4.

Attention, la prise en charge du type double est optionnelle et doit se conformer aux exigences de la norme IEEE-754 [13]. Tout code noyau qui utilise le type double ou ses dérivés vectoriels doit inclure la directive #pragma OPENCL EXTENSION cl_khr_fp64 : enable.

Une pseudo structure permet d'accéder au scalaire d'un vecteur ou d'en extraire de nouveaux vecteurs :

• accès à un scalaire par son nom jusqu'à la dimension 4 : utiliser les suffixes .x, .y, .z, .w ;
• accès à un scalaire par son rang : utiliser les suffixes .sN où N est un chiffre hexadécimal de 0 à F ;
• construction d'un vecteur à partir d'un autre vecteur par simple concaténation des suffixes (voir exemples à suivre) ;
• extraction de vecteur par moitié :

• • 1ère moitié (« gauche ») : suffixes .odd ou .lo,
  • 2ème moitié (« droite ») : suffixes .even ou .hi.

Les opérateurs habituels sont aussi disponibles pour les types vectoriels. Les deux vecteurs doivent être de même dimension :

• opération arithmétique (+, -, * , /, %) : application de l'opération sur les composantes deux à deux. Attention, dans le cas de la multiplication, il ne s'agit donc ni du produit matriciel, ni du produit vectoriel.
• comparaison (<, >, <=, >=, ==, !=) : comparaison des éléments deux à deux de deux vecteurs de même taille, donne un vecteur d'entier contenant le résultat de chaque comparaison.

Les mêmes opérateurs sont aussi applicables entre un scalaire et un vecteur. Dans ce cas, on obtient pour résultat un vecteur contenant le résultat de l'opération entre le scalaire et chaque composante du vecteur.

4.3 Synchronisation de l'hôte et des commandes

L'hôte communique avec les périphériques de calcul via des queues de commandes, c'est-à-dire des files d'attente de requêtes et d'événements. Comme déjà indiqué, la principale fonction à utiliser pour démarrer des tâches, clEnqueueNDRangeKernel, n'est pas bloquante et le programme de l'hôte continue à se dérouler en parallèle du noyau qui vient d'être lancé. Il en est de même pour la lecture et l'écriture de données de l'espace global par l'hôte. La fonction clEnqueueReadBuffer soumet la demande d'accès à la file d'attente de commandes, mais ne bloque pas forcément jusqu'à la disponibilité des données. Comme pour ce cas particulier, l'attente reste le cas d'usage le plus courant, le troisième paramètre est un booléen qui permet de préciser explicitement ce que l'on souhaite faire. Mais de façon générale, toutes les fonctions qui envoient des ordres dans la queue de commandes offrent les services suivants :

- définition des événements à attendre pour autoriser le traitement de la commande ;

- fourniture du descripteur de l'événement qui correspond à la terminaison de l'action demandée.

Pour compléter, la méthode clWaitForEvents permet à l'hôte d'attendre spécifiquement une commande.

Pour illustrer, imaginons quatre noyaux A, B1, B2 et C. B1 et B2 doivent chacun attendre la réalisation de A, et C celle des deux tâches B1 et B2. La synchronisation pourra être établie explicitement dans le code de l'hôte par clWaitForEvents ou mieux, quand aucun traitement intermédiaire par l'hôte n'est requis, directement lors de la programmation des tâches.

4.4 Synchronisation entre work-items

OpenCL propose un mécanisme de « barrière » pour organiser la collaboration des work-items d'un même work-group. Il repose sur la fonction barrier à invoquer depuis le noyau et qui reçoit en paramètre un champ de bits à définir à partir des deux constantes suivantes :

- CLK_LOCAL_MEM_FENCE : synchronisation de l'espace mémoire local.

- CLK_GLOBAL_MEM_FENCE : synchronisation de l'espace mémoire global.

La synchronisation d'un espace mémoire vise à s'assurer que toutes les modifications faites par les work-items dans cet espace sont bien propagées avant de poursuivre.

Une telle barrière permet de cadencer le déroulement de tous les work-items d'un work-group. C'est-à-dire que pour chaque work-item, le déroulement de processus bloque à l'encontre de la fonction barrier, jusqu'à ce que tous les autres work-items du même work-group l'aient atteinte. Il est donc nécessaire que tous les work-items la rencontrent. Dans le cas contraire, les work-items bloqués le seront au moins jusqu'à terminaison complète des autres work-items (ce comportement constaté sur mon système n'est pas spécifié par la norme). Donc il faudra être très rigoureux lors de l'utilisation de cette fonction dans des blocs conditionnels et dans des boucles :

- Toutes les branches conditionnelles doivent rencontrer le même nombre d'appels à barrier.

- Toutes les boucles qui la contiennent doivent effectuer exactement le même nombre d'itérations.

Avec ces précautions, les barrières offrent la possibilité de faire avancer les work-items ensemble, en quelque sorte pas à pas, et assurent la cohérences des espaces mémoire modifiés par chacun qui pourront alors être utilisés comme espaces partagés.

Examinons leur utilisation à l'aide d'un exemple. Supposons que l'on veuille normaliser un vecteur (diviser toutes ses composantes par son module, c'est-à-dire sa taille selon la distance euclidienne). Le calcul doit être divisé en deux grandes étapes : la détermination du module, puis son utilisation pour diviser chaque composante. Il faut évidemment attendre que le calcul du module soit terminé pour commencer les divisions et un seul work-item parmi l'ensemble sera chargé de calculer le module.

Ce premier exemple présente un problème et ne fonctionne que dans un cas particulier (qui est quand même acceptable dans beaucoup de situations) : on doit pouvoir confier l'intégralité du calcul à un seul work-group, c'est-à-dire qu'il faut que le vecteur traité soit de taille inférieure ou égale à MAX_GROUP_SIZE (ou MAX_WORK_ITEM_SIZE selon la dimension utilisée pour parcourir l'espace de travail). Contrairement à ce qu'une lecture trop rapide de la documentation peut laisser croire, CLK_GLOBAL_MEM_FENCE ne signifie pas que le verrou est global, il reste local à un work-group. Il n'y a aucun moyen de synchroniser des work-groups autre que de séquencer des noyaux, comme vu au paragraphe précédent. CLK_GLOBAL_MEM_FENCE assure uniquement une synchronisation de l'espace mémoire global avant de poursuivre, c'est-à-dire que toutes les altérations de l'espace mémoire global effectuées jusqu'à la barrière par chaque work-item du même work-group ont été prises en compte. L'utilisation d'une variable dans l'espace mémoire global n'est par conséquent pas d'une grande utilité dans le présent exemple. De plus, si des work-groups différents ne peuvent pas être synchronisés, il faut répéter le calcul du module pour chaque work-group :

Avec ces corrections, le calcul est correct quelles que soient les tailles relatives du vecteur et des work-groups. Il n'est cependant pas optimal puisque le calcul du module est intégralement répété pour chaque work-group. Une meilleure solution de ce point de vue est bien entendu de scinder le travail dans deux noyaux distincts qui seront exécutés l'un après l'autre, comme montré dans le paragraphe précédent. Il reste que pour de nombreux problèmes, la conservation en espace local ou global de résultats intermédiaires partagés entre work-items sera d'une grande aide pour rédiger des codes efficaces.

Conclusion

Nos machines actuelles, même les plus modestes, offrent aujourd'hui une puissance de calcul remarquable, mais encore faut-il y avoir accès. OpenCL le permet pour un composant sous-exploité, à part par quelques joueurs exigeants, c'est-à-dire le processeur graphique. A titre indicatif, le nettop Zotac utilisé pour l'élaboration de cet article s'est montré capable d'accomplir plus de 20 GFlops avec certains exemples joints au SDK CUDA. C'est remarquable pour une machine qui consomme moins de 30W et coûte moins de 300€ TTC. Mais la puissance de calcul brute ainsi mesurée n'est pas tout. La performance d'une application complète dépend aussi beaucoup de la bande passante entre le système hôte et le GPU, car elle est le passage obligé pour toutes les entrées/sorties des noyaux. Comme ce modèle est performant, surtout dans la situation où on répète le même calcul sur une masse importante de données, c'est un paramètre incontournable à prendre en considération.

De plus, OpenCL et ses concurrents, aussi bien CUDA que DirectCompute, ne tournent pas le dos au monde d'où ils sont originaires. Tous incluent des types de données spécialisés et des fonctions d'échange avec leur « ancêtre » graphique. C'est-à-dire qu’OpenCL permet en quelque sorte l'échange de pointeurs avec OpenGL pour travailler sur des espaces mémoire communs et, de ce fait, évite de solliciter l'hôte et son canal de communication avec le GPU quand les calculs portent directement sur des images.

Pour terminer, considérer OpenCL comme un outil dédié au GPGPU est assez réducteur. C'est d'abord une interface de programmation de ressources composites de calcul quelles qu'elles soient : toute architecture multicœur, homogène ou non, est visée. OpenCL apporte, par rapport à ses concurrents, une très grande portabilité vers toutes ces plates-formes, et de ce fait, une plus grande pérennité.

Enfin, pour être complet, il faut quand même signaler que même si OpenCL est un standard ouvert, toutes ses mises en œuvre disponibles aujourd'hui reposent sur une couche propriétaire (drivers + SDK et qui ne sont même pas distribués sous forme de paquets).

Références

[1] http://fr.wikipedia.org/wiki/3dfx

[2] The renderMan Interface, version 3.2.1, Pixar, November 2005

Part II, the RenderMan Shading Language, http://renderman.pixar.com/products/rispec/rispec_pdf/RISpec3_2.pdf

[3] http://www.nvidia.com/object/cuda_home_new.html

[4] http://www.khronos.org/registry/cl/

[5] NVidia CUDA - programming guide, version 3.0, 2/20/2010

http://developer.nvidia.com/object/cuda_3_0_downloads.html

[6] http://developer.amd.com/gpu/ATIStreamSDK/Pages/default.aspx

[7] ATI Stream SDK v2.01 - Installation notes, http://developer.amd.com/gpu/ATIStreamSDK/pages/Documentation.aspx

[8] The OpenCL specification, version 1.0, révision 0.48, Khronos OpenCL Working Group. 10/6/2009, http://www.khronos.org/registry/cl/specs/opencl-1.x-latest.pdf

[9] S. Devaux : « Initiation au calcul intensif sur PS3 », GNU Linux Magazine France hors-série n°43, pages 74 à 82, Août/Sept 2009

[10] http://llvm.org/

[11] http://www.alphaworks.ibm.com/tech/opencl

[12] http://fr.wikipedia.org/wiki/Scalable_Link_Interface

[13] http://fr.wikipedia.org/wiki/IEEE_754