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

Magazine
Marque
GNU/Linux Magazine
Numéro
129
Mois de parution
juillet 2010


Résumé
Un outil conçu pour dessiner, le processeur graphique, permet maintenant de réaliser des calculs de toutes sortes. La généralisation des shaders dans de nouveaux outils de développement comme OpenCL et CUDA permet à tout un chacun de disposer à des coûts très compétitifs d'une très grande puissance de calcul parallèle. Sans doute que l'histoire technologique n'est qu'une répétition comme l'Histoire en général : le dessin et la géométrie ont été, dans l'antiquité, le seul et probablement le premier système de calcul capable de déterminer avec une remarquable fiabilité pour l'époque, des grandeurs aussi diverses que la hauteur des pyramides ou la circonférence de la terre.

Body

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).

cpu-gpu

Architectures CPU et GPU

Pixar a introduit les shaders dans son outil RenderMan [2] pour permettre à ses artistes de décrire des processus d'ombrage ou de texture de façon souple à l'aide d'un langage de programmation qui leur est accessible. Ceci permet d'obtenir des surfaces d'aspect complexe avec un certain réalisme à partir d'éléments descriptifs relativement simples qui sont itérés par le moteur de rendu lors du calcul des images. Dans les première cartes graphiques 3D, la section de rendu était relativement rigide, seuls quelques paramètres étaient disponibles pour agir sur les processus d'ombrage et de lissage. OpenGL et DirectX ont tous les deux par la suite introduit des shaders pour pouvoir programmer spécifiquement ces processus. Plus tard, de nouveaux shaders ont été introduits pour agir directement sur le modèle 3D et non plus uniquement sur les pixels. Il est alors devenu très tentant d'utiliser ces nouvelles facultés de calcul vectoriel à d'autres fins que le rendu graphique.

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.

/* ### Addition de vecteurs avec C for CUDA ### */

#include <stdio.h>

// definition d'une fonction CUDA

// prise en charge par le GPU.

__global__ void addVect(float* in1, float* in2, float* out) {

 int i = threadIdx.x;

 out[i] = in1[i] + in2[i];

}

// Code C "normal", est exécuté par le processeur de l'hôte.

int main() {

    // données d'entrée

    float v1[]={1, 2, 3, 4, 5, 6, 7, 8, 9, 10};

    float v2[]={10.9, 11.8, 12.7, 13.6, 14.5, 15.4, 16.3, 17.2, 18.1, 19.0};

    int memsize = sizeof(v1);

    int vsize = memsize/sizeof(float);

// données de sortie

    float res[vsize];

    // allocation mémoire sur le GPU

    float * Cv1;

    cudaMalloc((void **)&Cv1, memsize);

    float * Cv2;

    cudaMalloc((void **)&Cv2, memsize);

    float * Cres;

    cudaMalloc((void **)&Cres, memsize);

// copie données d'entrée sur le GPU

    cudaMemcpy(Cv1, v1, memsize, cudaMemcpyHostToDevice);

    cudaMemcpy(Cv2, v2, memsize, cudaMemcpyHostToDevice);

    // déclanche l'execution de la fonction CUDA sur le GPU

    // grille de 1 bloc de même taille que le vecteur

    addVect<<<1,vsize>>>(Cv1, Cv2, Cres);

// copie données de sortie du GPU

    cudaMemcpy(res, Cres, memsize, cudaMemcpyDeviceToHost);

    // affichage resultat

    int i=0;

    printf("res = { ");

    for (i=0;i<vsize;i++) {

     printf("%f ",res[i]);

}

    printf("}\n");

}

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>>>.

/* ### Compilation et exécution d'un programme C for CUDA ### */

$ nvcc --link -o addvec addvect.cu

$ ./addvec

res = { 11.900000 13.800000 15.700000 17.600000 19.500000 21.400000 23.299999 25.200001 27.100000 29.000000 }

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 :

/* ### Initialiser premier périphérique OpenCL ### */

#include <CL/cl.h>

/* ... */

 cl_platform_id oclPF;

 cl_device_id oclDev[1]; /* les fonction API OpenCL travaillent sur des tableaux de cl_device_id */

 cl_int oclErr;

 oclErr=clGetPlatformIDs(1, &oclPF, NULL);

 if (oclErr != CL_SUCCESS) {

  /* gestion erreur */

  /* ... */

 }

 oclErr=clGetDeviceIDs(oclPF, CL_DEVICE_TYPE_ALL, 1, oclDev, NULL);

 if (oclErr!=CL_SUCCESS) {

/* gestion erreur */

  /* ... */

 }

 cl_context oclContext=clCreateContext(0, 1, oclDev, NULL, NULL, &oclErr);

 if (oclErr!=CL_SUCCESS) {

  /* gestion erreur */

  /* ... */

 }

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.

/* ### définition et compilation du code OpenCL ### */

const char *oclSource=

    "__kernel void vectAdd(__global const float* a,\n"

    "        __global const float *b,\n"

    "        __global float * res) {\n"

    "    int i=get_global_id(0);\n"

    "    res[i]=a[i]+b[i];\n"

    "}\n";

cl_program oclPgm = clCreateProgramWithSource(oclContext, 1, (const char **)&oclSource, NULL, &oclErr);

oclErr = clBuildProgram(oclPgm, 0, NULL, NULL, NULL, NULL);

cl_command_queue oclCmdQueue = clCreateCommandQueue(oclContext, oclDev[0], 0, &oclErr);

cl_kernel oclKernel = clCreateKernel(oclPgm, "vectAdd", &oclErr);

code

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.

/* ### Initialisation donnée et mémoire ### */

/* données dans l'espace mémoire du système hôte */

#define VSIZE 512

float in1[VSIZE];

float in2[VSIZE];

float out[VSIZE];

/* ... */

/* allocation mémoire du périphérique OpenCL et initialisation avec la mémoire de l'hôte */

cl_mem oclIn1=clCreateBuffer(oclContext, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, VSIZE*sizeof(cl_float), &in1,&oclErr);

cl_mem oclIn2=clCreateBuffer(oclContext, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, VSIZE*sizeof(cl_float), &in2,&oclErr);

cl_mem oclOut=clCreateBuffer(oclContext, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR, VSIZE*sizeof(cl_float), &out,&oclErr);

code

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.

/* ### Invocation d'un noyau OpenCL ### */

/* association des tampons aux paramètres de la fontion */

clSetKernelArg(oclKernel, 0, sizeof(cl_mem), (void *)&oclIn1);

clSetKernelArg(oclKernel, 1, sizeof(cl_mem), (void *)&oclIn2);

clSetKernelArg(oclKernel, 2, sizeof(cl_mem), (void *)&oclOut);

/* lancement du calcul */

int arraySize=VSIZE;

clEnqueueNDRangeKernel(oclCmdQueue, oclKernel, 1, 0, &arraySize, 0, 0, 0, 0);

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.

/* ### obtention des resultats et libération des buffers ### */

/* copie mémoire périphérique dans mémoire principale */

clEnqueueReadBuffer(oclCmdQueue, oclOut, CL_TRUE, 0, VSIZE*sizeof(cl_float), &out, 0, 0, 0);

/* libération des tampons du périphériques */

clReleaseMemObject(oclIn1); 

clReleaseMemObject(oclIn2); 

clReleaseMemObject(oclOut); 

/* afficher quelques valeurs */

printf("%f op %f = %f\n", in1[12], in2[12], out[12]);

printf("%f op %f = %f\n", in1[128], in2[128], out[128]);

printf("%f op %f = %f\n", in1[327], in2[327], out[327]);

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.

/* ### compilation et execution ### */

sdevaux@arro:src$ gcc -o cladd.obj -c cladd.c

sdevaux@arro:src$ gcc -o cladd cladd.obj -lOpenCL

sdevaux@arro:src$ ./cladd

12.000000 op 12000.000000 = 12012.000000

128.000000 op 128000.000000 = 128128.000000

327.000000 op 327000.000000 = 327327.000000

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 :

/* ### Inventaire Atom/ION ### */

0: NVIDIA CUDA

  0/0: ION

     MAX_WORK_GROUP_SIZE:     512

     MAX_WORK_ITEM_DIMENSIONS: 3

     MAX_WORK_ITEM_SIZES:      [ 512 512 64 ]

/* ### Inventaire Core/ATI ### */

0: ATI Stream

0/0: Intel(R) Core(TM)2 Duo CPU     E6750 @ 2.66GHz

    MAX_WORK_GROUP_SIZE:     1024

    MAX_WORK_ITEM_DIMENSIONS: 3

    MAX_WORK_ITEM_SIZES:      [ 1024 1024 1024 ]

0/1: ATI RV770

    MAX_WORK_GROUP_SIZE:     256

    MAX_WORK_ITEM_DIMENSIONS: 3

    MAX_WORK_ITEM_SIZES:      [ 256 256 256 ]

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.

La mise en œuvre OpenCL sur le GPU par le SDK ATI semble s'appuyer sur les bibliothèques X. En effet, l'inventaire des périphériques montré plus haut n'est possible sur le système ATI que depuis une session graphique (ou au moins l'accès à un DISPLAY sur une session ouverte), sans cela le GPU n'est pas détecté et le message « No protocol specified » apparaît. Pour CUDA, il n'y a en revanche aucun besoin de serveur X actif tant qu'on ne fait aucun affichage graphique.

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.

/* ### Extrait code compilé par ION ### */

//

// Generated by NVIDIA NVPTX Backend for LLVM

//

.version 1.5

.target sm_11, texmode_independent, map_f64_to_f32

// Global Launch Offsets

.const[0] .s32 %_global_block_offset[3];

.const[0] .s32 %_global_launch_offset[3];

.const[0] .s32 %_global_num_groups[3];

.const[0] .s32 %_global_size[3];

.const[0] .u32 %_work_dim;

- 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.

/* ### Code compilé pour périphérique Intel Core ### */

file:///tmp/OCLwZCwsi.so

/* ### Examen du fichier cible du descripteur ATI / Intel Core ### */

sdevaux@levia-desktop:~/src/addvect$ nm /tmp/OCLwZCwsi.so

00000fb0 r .str

00000fb2 r .str1

00000fb4 r .str2

000020a0 D CubeRootTable_

000028a0 D FloatReciprocalTable_

00002000 a _DYNAMIC

00002078 a _GLOBAL_OFFSET_TABLE_

00000620 T __OpenCL_vectAdd_kernel

00003130 D __OpenCL_vectAdd_metadata

00000650 T __OpenCL_vectAdd_stub

00000680 T __OpenCL_vectMul_kernel

000031a4 D __OpenCL_vectMul_metadata

000006b0 T __OpenCL_vectMul_stub

00000a30 T __amdrt_cvt_f32_to_u64

00000b40 T __amdrt_cvt_f64_to_u64

00000c50 T __amdrt_div_i64

00000710 T __amdrt_div_u64

00000df0 T __amdrt_mod_i64

00000860 T __amdrt_mod_u64

00003218 A __bss_start

00003218 A _edata

0000321c A _end

[...]

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

- work-group : groupe de travail ;

- work-item : unité de travail ;

- compute-device : périphérique de calcul, système de calcul construit sur une architecture particulière. Il peut s'agir d'une carte multiprocesseur qui embarque sa propre mémoire, ou une carte graphique pouvant contenir une ou plusieurs puces GPU ;

- compute-unit : unité de calcul assimilable à un processeur ;

- processing-element : élément de calcul, cœur ou unité arithmétique et logique (UAL) d'un cœur.

L'association (entité de calcul OpenCL/élément matériel présentés ci-dessus) est donnée à titre indicatif pour illustration. La réalité peut être beaucoup plus complexe puisque les notions d'UAL ou de cœurs ne sont pas strictement équivalentes entre un processeur conventionnel (x86, multicœur ou non), un GPU ou une architecture hybride (Cell BE [11]). A charge du driver OpenCL de chaque architecture d'affecter les structures matérielles aux structures logiques équivalentes.

Dans le modèle CUDA, des notions équivalentes existent, puisqu'on y trouve également une hiérarchie à deux niveaux : grille et blocs. Et bien sûr, CUDA est aussi capable de prendre en charge un système composite incluant plusieurs cartes graphiques (cartes NVidia obligatoirement et de préférence avec le même modèle de GPU), interconnectées par SLI [12].

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.

/* ### Itération séquentiel à 3 dimension. ### */

int i,j,k;

for (i=0;i<X_MAX;i++) {

 for (j=0;i<Y_MAX;j++) {

  for (k=0;k<Z_MAX;k++) {

/* calcul : traitement d'un élément i,j,k */

   out[i][j][k] = calcelem(i,j,k,in);

}

 }

}

/* ### Noyau OpenCL à trois dimensions ### */

float calcelem(float i, float j, float k, __global float * in) {

float res;

    /* faire ici le calcul qui va bien avec in[i][j][k] */

 return res;

}

__kernel void func(__global float* in, __global float* _out) {

 int i = get_global_id(0);

 int j = get_global_id(1);

 int k = get_global_id(2);

 out[i][j][k] = calcelem(i,j,k,in);

}

/* ### Invocation du noyau ### */

size_t dimensions[3] = {X_MAX, Y_MAX, Z_MAX};

clEnqueueNDRangeKernel(oclCmdQueue, oclKernel, 3, 0, dimensions, 0, 0, 0, 0);

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.

Si on se base par exemple sur les capacités de la plate-forme ION, que nous avons découvertes plus haut : un work-group contient au plus 512 work-items, sur trois dimensions avec 512 max pour les deux premières et 64 pour la dernière. Examinons l'applicabilité de quelques exemples de définition de work-groups :

- {500} : valide ;

- {1024} : invalide car dépasse à la fois taille maximale absolue d'un work-group et taille maximale de la première dimension ;

- {70, 2, 2} : valide, 280 work-items ;

- {2,2,70} : invalide, 280 work-items seulement, mais limite de la troisième dimension dépassée ;

- {2,2,2,2} : invalide, 16 work-items, mais nombre de dimensions dépassé.

La taille totale des données à traiter a également son importance puisqu'il faudra pouvoir découper cette grandeur en lots de la taille des work-groups. Comme d'habitude dans notre monde informatique binaire, si on a le choix, travailler avec une puissance de deux facilitera grandement le découpage et garantira un bon rendement.

- 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.

/* ### Vecteurs 4D ### */

float4 a = (float4)(1.0f, 2.0f, 3.0f, 4.0f);

float4 b = (float4)(0.4f, 0.3f, 0.2f, 0.1f);

float4 c = a + b;

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.

/* ### Utilisation des type vectoriels ### */

int4 p=(1,2,3,4);

int x = p.x;      // x = 1

int y = p.y;      // y = 2

int4 pt = p.xyzw; // pt = (1,2,3,4)

int4 tp = p.wzyx; // tp = (4,3,2,1)

int4 pp = p.xyzw + p.zyxw; // pp = (4,4,4,8);

int4 dp = 2 * p; // dp = (2,4,6,8)

int4 p2 = p * p; // p2 = (1,4,9,16)

int4 cp1 = p.xyzw < p.yyyy; // cp1=(-1,0,0,0)

int4 cp2 = p.xyzw <= p.yyyy; // cp2=(-1,-1,0,0)

int16 v = (1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);

int s0 = v.s0;    // s0 = 1;

int s15 = v.sF;   // sF = 16;

int2 c = v.s9A;   // c = (10,11)

int mv = v.lo;    // mv = (1,2,3,4,5,6,7,8)

int qv = v.hi.lo // qv = (9,10,11,12)

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.

/* ### Attente interceptée par l'hôte ### */

cl_kernel kA, kB1, kB2, kC;

cl_commend_queue oclCvdQueue;

/*

* Initialisation divers à placer ici (noyaux, queue, arguments, dimensions, etc)

*/

cl_event finA[1];

cl_event finB[2];

clEnqueueNDRangeKernel(oclCmdQueue, kA, dims, 0, dimSpec, 0, 0, finA);

clWaitForEvent(1,finA); /* bloquer ici jusqu'à fin de kA */

clEnqueueNDRangeKernel(oclCmdQueue, kB1, dims, 0, dimSpec, 0, 0, &(finB[0]));

clEnqueueNDRangeKernel(oclCmdQueue, kB2, dims, 0, dimSpec, 0, 0, &(finB[1]));

clWaitForEvent(2,finB); /* bloquer ici jusqu'à fin de kB1 et kB2 */

clEnqueueNDRangeKernel(oclCmdQueue, kC, dims, 0, dimSpec, 0, 0, 0);

/* ### Soumission des travaux avec dépendances ### */

cl_kernel kA, kB1, kB2, kC;

cl_commend_queue oclCvdQueue;

/*

* Initialisation divers à placer ici (noyaux, queue, arguments, dimensions, etc)

*/

cl_event finA[1];

cl_event finB[2];

clEnqueueNDRangeKernel(oclCmdQueue, kA, dims, 0, dimSpec, 0, 0, finA);

clEnqueueNDRangeKernel(oclCmdQueue, kB1, dims, 0, dimSpec, 1, finA,&(finB[0]));

clEnqueueNDRangeKernel(oclCmdQueue, kB2, dims, 0, dimSpec, 1, finA,&(finB[1]));

clEnqueueNDRangeKernel(oclCmdQueue, kC, dims, 0, dimSpec, 2, finB, 0);

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.

/* ### noyau normalisation 1 ### */

__global float module;

__kernel void normvect(__global const float* in,

        __global float* out) {

    int i = get_global_id(0);

    if (i == 0) {

        /* le premier work-item calcule le module. */

        int size = get_global_size(0);

        module = 0;

        int j;

  for (j=0; j<size; j++) {

            module += in[j] * in[j];

     }

        module = sqrt(module);

}

    /* attendre fin du calcul du module */

    barrier(CLK_GLOBAL_MEM_FENCE);

    out[i] = in[i] / 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 :

/* ### noyau normalisation 2 ### */

__local float module;

__kernel void normvect(__global const float* in,

        __global float* out) {

    int i=get_global_id(0);

 if (get_local_id(0) == 0) {

/* le premier work-item du groupe calcule le module. */

        int size = get_global_size(0);

        module = 0;

        int j;

  for (j=0; j<size; j++) {

            module += in[j] * in[j];

     }

        module = sqrt(module);

}

    /* attendre fin du calcul du module */

    barrier(CLK_LOCAL_MEM_FENCE);

    out[i]= in[i] / module; 

}

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