Une introduction à CUDA

Thibaut Cuvelier

Une introduction à CUDA et au calcul sur GPU, comparativement avec les CPU.
Avant la fin, vous pourrez écrire vos premiers kernels.
Cette introduction se base sur CUDA 2.1 et 2.2.

N'hésitez pas à commenter cet article !
Commentez

- Titre : Une introduction à CUDA

- Auteur : Thibaut Cuvelier

- Parution : 4 avril 2009

- Mise à jour : 2 février 2020

- Licence : Copyright ® 2009 Thibaut Cuvelier. Aucune reproduction, même partielle, ne peut être faite de ce site et de l'ensemble de son contenu : textes, documents, images, etc. sans l'autorisation expresse de l'auteur. Sinon vous encourez selon la loi jusqu'à trois ans de prison et jusqu'à 300 000 € de dommages et intérêts.

Sommaire

- I. GPGPU

- II. CUDA

- II-A. Pilote

- II-B. Runtime

- II-C. Bibliothèques

- III. Un peu de vocabulaire

- IV. CPU et GPU

- IV-A. Survol de quelques différences

- IV-B. Précision des calculs

- IV-C. GPU

- IV-C-1. Mémoires

- IV-C-1-a. Mémoire globale

- IV-C-1-b. Mémoire locale

- IV-C-1-c. Mémoire constante

- IV-C-1-d. Mémoire des textures

- IV-C-1-e. Mémoire partagée

- IV-C-1-f. Registres

- IV-C-1-g. Mémoire système

- IV-C-2. Shaders

- IV-C-2-a. Plus de précisions

- IV-C-2-b. Limites

- IV-D. CPU

- IV-D-1. Mémoire cache

- IV-D-2. Pipelines d'instructions

- IV-D-3. Exécution superscalaire

- V. Les mains dans le cambouis

- V-A. Les kernels

- V-B. Qualifieurs de kernels

- V-B-1. __global__

- V-B-2. __device__

- V-B-3. __host__

- V-C. Configuration de l'exécution

- V-C-1. Dg

- V-C-2. Db

- V-C-3. Ns

- V-C-4. S

- V-D. Qualificateurs de variables

- V-D-1. __device__

- V-D-2. __constant__

- V-D-3. __shared__

- V-D-4. Généralités

- V-E. Compilation

- VI. Le modèle de programmation

- VI-A. Parallélisme des données

- VI-B. Structure du programme

- VI-C. L'exemple : la multiplication de matrices carrées

- VI-D. Séparer les opérations à effectuer

- VII. Conclusions

- VII-A. Et chez AMD/ATI ?

- VII-B. Intégration à Visual Studio

- VII-C. Avec d'autres environnements de compilation

- VII-D. Déploiement

- VII-E. Quels fichiers inclure ?

- VIII. Divers

- VIII-A. Références

- VIII-B. Voir aussi

- VIII-B-1. GPGPU

- VIII-B-2. CUDA

- VIII-B-3. Bibliothèques

- VIII-B-4. Matériel

- VIII-B-5. Autres

- VIII-C. Remerciements

I - GPGPU

Les constructeurs ont décidé de créer des langages qui permettent d'exploiter les possibilités de ces processeurs graphiques. Ils n'ont pas été les seuls.

Par exemple, l'université de Stanford a créé le BrookGPU, le tout premier langage, un dérivé du C, qui permet d'utiliser les API DirectX et OpenGL, ainsi que GLSL ou CG. L'avantage de ces solutions est qu'elles sont utilisables sur tous les GPU qui supportent DirectX et/ou OpenGL, c'est-à-dire la plus grande majorité d'entre eux, et la totalité ces dernières années. Cependant, cette universalité se traduit aussi par un manque de performances par rapport à d'autres bibliothèques plus proches du matériel.

Ainsi, ATI a développé Close to Metal, une bibliothèque très bas niveau. Cette bibliothèque sera suppléée par Stream, mais cette dernière est plus spécifiquement dirigée vers les processeurs FireStream, prévus pour le calcul.

Ensuite vient NVIDIA, avec CUDA, une technologie disponible sur toutes les cartes graphiques grand public depuis la série des GeForce 8000 et sur tous les supercalculateurs Tesla.

II - CUDA

Ou Compute Unified Device Architecture.

C'est la réponse de NVIDIA aux demandes sans cesse croissantes de puissance de calcul. Cette bibliothèque, dévoilée en 2007, permet d'employer la puissance de calcul des GPU. Elle n'est que la partie logicielle du tout : il faut encore une carte graphique compatible.

CUDA supporte plusieurs langages : le C, le C + + et le Fortran. Vous pouvez donc utiliser conjointement ces trois langages dans vos fonctions et vos kernels.

Il existe déjà quelques wrapper pour CUDA : PyCUDA, destiné à Python, ainsi que JCublas, JCufft et JCudpp, sans oublier CUDA lui-même, avec jCUDA pour Java, sans oublier CuBLAS.Net, un wrapper de CuBLAS pour le CLR .Net.

CUDA est constitué d'un pilote, déjà intégré aux ForceWare les plus récents ; d'un runtime ; et de quelques bibliothèques. CUDA est aussi un langage, dérivé du C (mais n'apportant que peu de modifications : 9 nouveaux mots-clés, 24 nouveaux types et 62 nouvelles fonctions). Ces extensions nécessitent leur compilateur, lui aussi fourni.

CUDA est prévu pour s'exécuter sur un GPU, mais il est aussi disponible sur CPU, en émulation. Les performances sont alors bien moindres, mais cela peut être utile pour tester ses applications sans GPU compatible.

L'API CUDA est de haut niveau : vous ne vous occupez donc pas du GPU directement. CUDA en est une couche d'abstraction.

Voici, graphiquement représentées, toutes les composantes de CUDA et de son utilisation.

II-A - Pilote

- Rôle : transmettre les calculs de l'application au GPU ;

- Distribution avec les ForceWare 178.08 et plus récents ;

- Inconvénient : pas d'automatismes.

II-B - Runtime

- Rôle : interface entre le GPU et l'application, en fournissant quelques automatismes ;

- Distribution : en même temps que le pilote ;

- Inconvénient : impossibilité d'optimiser à partir d'un certain point.

II-C - Bibliothèques

Pour le moment, CUDA est livré avec CuBLAS et CuFFt, respectivement les implémentations de BLAS (une bibliothèque d'algèbre) et de la transformation rapide de Fourier (utilisée en analyse de Fourier et en traitement du signal). La dernière n'est pas inspirée d'une bibliothèque préexistante.

Ces implémentations reprennent le fonctionnement des bibliothèques originelles (CuBLAS), ou bien des algorithmes les plus performants (CuFFT) et les optimisent au maximum pour CUDA.

III - Un peu de vocabulaire

Nous allons continuer cette introduction avec un peu de vocabulaire inhérent à la programmation avec CUDA.

L'hôte est le CPU, c'est lui qui demande au périphérique (le GPU) d'effectuer les calculs.

Un kernel est une portion parallèle de code à exécuter sur le périphérique. Chacune de ses instances s'appelle un thread.

Une grille est constituée de blocs. Chaque bloc est constitué de threads.

Un bloc est un élément des calculs, dissociable d'autres blocs : les blocs ne doivent donc pas être exécutés dans un certain ordre : parallèlement, consécutivement ou toute autre combinaison est possible. C'est pourquoi les threads ne peuvent communiquer qu'avec des threads du même bloc.

Un warp est un ensemble de 32 threads, envoyés ensemble à l'exécution et exécutés simultanément. Quel que soit le GPU utilisé, quelle que soit la quantité de données à traiter, dans n'importe quel cas, un warp sera exécuté sur deux cycles. On peut être sûr et certain qu'ils le seront. Ceci pourra vous aider lors de la conception de vos algorithmes. Par exemple, Mark Harris, chercheur pour NVIDIA dans le rendu graphique en temps réel, fondateur du site GPGPU, utilise cette donnée pour dérouler ses boucles.

Un petit parallèle avec le matériel. Un thread est exécuté par un processeur : posons donc l'égalité entre le thread et le processeur. Ainsi, le bloc est le multiprocesseur, tandis que la grille représente l'entièreté de la carte.

Le calcul hétérogène est l'utilisation des deux types de processeurs disponibles sur nos ordinateurs : les CPU et les GPU. Il s'agit donc d'utiliser le bon type de processeur pour la bonne tâche.

Vous voici prêt pour partir à l'attaque !

IV - CPU et GPU

IV-A - Survol de quelques différences

La puissance de nos GPU n'a de cesse d'augmenter depuis quelques années. À un point qu'il est désormais possible de les utiliser pour réaliser des calculs autres que pour des jeux. En effet, parmi les CPU, un Intel Pentium 4 cadencé à 3 GHz fournit 4,8 GFlops, un Intel Core 2 Duo E6750 (2,66 GHz), 14,2 GFlops ; chez les GPU, on change de catégorie : la GeForce 9800 GTX, 420 GFlops, pour 675 MHz seulement.

Cependant, ces différences énormes s'expliquent très facilement, explications dans ce tableau.

CPU (hors SIMD)

GPU

Nombre de tâches

Une seule et unique

Le plus grand nombre

Variété des tâches

Toutes possibles

Restreinte

Subdivision de la tâche

Aucune : tout en un coup

Maximale, pour mieux la répartir sur les différentes unités de calcul

Il ne faut pas oublier de préciser que les GPU préfèrent travailler avec des vecteurs. Dans le cas contraire, les gains sont réellement minimes.

Les deux types de processeurs travaillent de façon radicalement différente. L'emploi de GPU à la place de CPU ne se fait donc pas en un tour de main : il faut repenser le calcul pour l'adapter au type de processeurs désiré. Si l'on ne change pas sa manière de penser, autant continuer de produire son électricité à la pomme de terre, qui permet quand même de produire assez pour éclairer quelques centimètres ; tandis que la centrale électrique permet d'éclairer des villes entières.

Pour le grand public, les prix se tiennent : un E6750 coûte, actuellement, 140 € ; une 9800 GTX, 150 €. Leurs éditions professionnelles sont légèrement différentes : 1500 $ pour un NVIDIA Tesla S870 plafonnant à 2 Tflops, contre 200 000 $ pour un IBM BlueGene de même puissance. Ici, on remarque bien l'un des grands avantages du GPGPU.

On peut considérer des racks de cartes Tesla comme des supercalculateurs. En effet, ce sont eux qui calculent. Cependant, un ou plusieurs CPU les orchestrent, en plus de leur donner la masse de travail.

Aussi, les GPU ont été, à la base, destinés à et spécialisés pour des calculs intensifs. Ceci leur permet de réserver plus de transistors au traitement des données, au lieu de les utiliser pour le cache ou bien pour la gestion des flux d'entrée ou de sortie.

Ainsi, un GPU doit être constitué de beaucoup de processeurs pour ces calculs : un GPU comporte au strict minimum 32 processeurs (240 pour le T10, 128 en moyenne) et ce, depuis plus qu'un temps certain. Ces processeurs sont les équivalents des cœurs de nos CPU, qui en comportent, en moyenne, deux depuis quelques années et, dans les années à venir, 80. Nous sommes donc bien loin des GPU !

IV-B - Précision des calculs

Les GPU actuels, avec CUDA, n'ont qu'une précision FP32, sur 32 bits. Il faut se tourner vers les solutions d'ATI/AMD pour une précision double sur 64 bits, ou bien vers des GPU plus chers, comme les Tesla ou les Quadro, ou bien récents, comme tous les GPU basés sur le GT200 (GeForce GTX260 à GTX295).

Tous les processeurs ne fonctionnent pas à la même précision : sur les premières GeForce compatibles CUDA, tous sont FP32. Sur un T10, huit unités sont FP32 et une seule FP64. Chez AMD, pour huit unités FP64, il y a quatre unités FP32.

Le peu d'unités dédiées au calcul à double précision sur les Tesla et autres explique leur faible puissance à ce niveau de précision, en comparaison de la simple précision ou bien des solutions d'AMD. Ainsi, pour du calcul en haute précision, les solutions NVIDIA tous publics ne sont pas encore au point (AMD ne propose plus de GPGPU pour la même gamme).

Actuellement, tous les processeurs supportent la double précision sur 64 bits.

Plus précisément, NVIDIA met à disposition la liste des écarts avec les standards, ainsi que ses limitations.

1. Les additions et soustractions sont souvent associées en une seule instruction ;

2. La division et la racine carrée sont implémentées par la réciproque, non conformément aux standards ;

3. Pour la multiplication et l'addition, il n'est possible que d'arrondir vers le nombre pair le plus proche ;

4. Il n'y a pas de possibilité d'arrondi configurable dynamiquement ;

5. Il n'y a pas de signalisation de NaN (Not a Number) ;

6. Il n'y a pas de mécanisme de détection d'exception, qui sera masquée selon les standards ;

7. Les opérandes de source dénormalisée tendent vers 0 ;

8. Le résultat d'une opération avec NaN est un NaN canonique de la forme 0x7fffffff ;

9. En accord avec les standards, si un NaN est passé à min() ou à max(), l'autre sera retourné.

IV-C - GPU

IV-C-1 - Mémoires

IV-C-1-a - Mémoire globale

CUDA est capable de lire et d'écrire sur la mémoire embarquée dans la carte graphique. Ces opérations portent, respectivement, les doux noms de gathering et de scattering.

La mémoire globale est la mémoire utilisable de n'importe quel endroit de CUDA, avec les mêmes performances à la clé : cette mémoire n'est pas cachée et il faut attendre 400 à 600 cycles avant d'y accéder. Ce qui laisse un multiprocesseur inactif pendant ce temps.

Pourquoi une telle latence ?

La mémoire globale est, en général (dans tous les cas, jusqu'à présent), de la DRAM.

Cette mémoire est très bon marché : 1,50 $ en septembre 2008, pour les intégrateurs ! Ceci lui permet d'être utilisée comme mémoire principale de nos ordinateurs.

De plus, elle se révèle compacte : on en fait tenir des Go sans problème sur des cartes !

Pourtant, cette mémoire a un problème et il s'agit de la latence. Elle monte sans problème jusqu'à 30 ns, ce qui représente quand même déjà 30 cycles ! Et sans compter les bus entre le multiprocesseur et la mémoire.

Finalement, cette mémoire n'est pas cachée.

IV-C-1-b - Mémoire locale

Cette mémoire est, à l'instar de la mémoire globale, non cachée et avec une latence très élevée.

Cette mémoire n'est utilisée que pour certaines variables, qui y sont placées automatiquement. En effet, certains tableaux, normalement placés dans les registres, sont trop grands : il leur faut donc un espace plus grand, qu'offre la mémoire locale.

IV-C-1-c - Mémoire constante

La mémoire constante est cachée : la lecture depuis cette mémoire ne coûte qu'un cycle. Pour tous les threads d'un demi-warp, la lecture depuis la mémoire constante est aussi rapide que depuis un registre, aussi longtemps que tous les threads lisent le même emplacement mémoire. Le coût de lecture augmente linéairement avec le nombre d'adresses différentes demandées par les threads. Il est recommandé que tous les threads d'un warp utilisent la même adresse et non seulement ceux de demi-warps, vu que les périphériques futurs le requerront pour un fonctionnement optimal.

Chaque multiprocesseur dispose d'une mémoire réservée aux constantes, d'une taille de 8 ko, dans le cas des GeForce 8800.

IV-C-1-d - Mémoire des textures

Cet espace mémoire est caché, le coût de la lecture est donc très faible.

Cette mémoire est optimisée pour un espace à deux dimensions, ainsi, les threads d'un même warp qui lisent à des adresses proches auront des performances optimales.

Aussi, elle est prévue pour des demandes de flux avec une latence constante.

La lecture des mémoires du périphérique par le mécanisme des textures peut être un avantageux substitut à la lecture depuis les mémoires globale ou constante.

Les textures seront approfondies plus tard, mais voici un avant-goût.

Les textures permettent vraiment de simplifier le traitement d'images : elles permettent la mise en œuvre de filtrages bilinéaires et trilinéaires très facilement et l'accès aléatoire aisé aux pixels.

IV-C-1-e - Mémoire partagée

Cette mémoire est présente sur le chipset, ce qui lui permet d'être assez rapide, plus que la mémoire locale.
En fait, pour tous les threads d'un warp, accéder à cette mémoire est aussi rapide que d'accéder à un registre, tant qu'il n'y a pas de conflit entre les threads.

Pour permettre une bande passante assez élevée, la mémoire partagée est divisée en modules de mémoire, les banques, qui peuvent être accédées simultanément. Ainsi, n lectures ou écritures qui tombent dans des banques différentes peuvent être exécutées simultanément dans un warp, ce qui permet d'augmenter sensiblement la bande passante, qui devient n fois plus élevée que celle d'un module.

Cependant, si deux demandes tombent dans la même banque, il y a un conflit de banques et l'accès doit être sérialisé. Le matériel divise ces requêtes problématiques en autant de requêtes que nécessaire pour qu'aucun problème n'ait lieu, ce qui diminue la bande passante d'un facteur équivalent au nombre de requêtes total à effectuer.

Pour des performances maximales, il est donc très important de comprendre comment les adresses mémoires sont reliées aux banques, pour pouvoir prévoir les requêtes et, ainsi, minimiser les conflits.

Dans le cas d'un espace en mémoire partagée, les banques sont organisées pour que des mots successifs de 32 bits soient assignés à des banques successives. Chaque banque a une bande passante de 32 bits tous les deux cycles d'horloge.

Pour le moment, un warp a une taille de 32 threads et il y a 16 banques.

Une requête en mémoire partagée pour un warp est divisée en deux : une partie pour le premier demi-warp, une autre, pour l'autre moitié. Ce qui a pour conséquence qu'il ne peut y avoir de conflit entre chaque demi-warp. Les conflits seront détaillés plus tard.

Actuellement, la mémoire partagée atteint un total de 16 ko, 1 ko pour chaque banque.

Pour résumer ceci, voici un schéma qui reprend l'essentiel des caractéristiques présentées ici.

IV-C-1-f - Registres

Généralement, l'accès à un registre ne prend pas un seul cycle supplémentaire par instruction, mais des retards peuvent apparaître, suite aux dépendances de lecture après écriture et des conflits qui peuvent se produire.

Les retards introduits pas les dépendances peuvent être ignorés, dès qu'il y a au moins 192 threads actifs par multiprocesseur, qui permettent de les cacher.

Le compilateur et l'organisateur des threads organisent les instructions pour des performances optimales, qui nécessitent d'éviter les conflits avec les banques. Le meilleur moyen d'obtenir de bonnes performances est d'utiliser un multiple de 64 comme nombre de threads par bloc. Une application n'a strictement aucun moyen de contrôler ces conflits.

Chaque multiprocesseur dispose de 8192 registres.

IV-C-1-g - Mémoire système

Depuis les GT200 (GeForce GTX 260 à 295), il est désormais possible d'utiliser la mémoire principale du système, alias RAM, grâce à CUDA 2.2.

Les appels à cette mémoire ne peuvent être fréquents : ils sont encore plus lents que les appels à la mémoire locale (700 à 800 cycles de latence !). Mais la RAM est disponible, de nos jours, en quantités plus grandes que celle disponible sur nos GPU.

IV-C-2 - Shaders

Les calculs demandés à CUDA sont, pour le moment, effectués sur les unités de shaders, les processeurs les plus rapides sur les GPU. Par exemple, les GeForce 8800 GTX ont des unités cadencées à 1,2 GHz.

Chaque unité de traitement des shaders est, comme montré ci-dessus, constituée de Texture Processor Clusters (TPC).

Chacun de ces clusters est fait d'une unité de traitement des textures (TEX) et de deux unités de traitement des flux (SM, Streaming Multiprocessor).

Vous n'avez pas vraiment besoin d'en savoir beaucoup plus pour pouvoir aborder CUDA. Cependant, si vous en voulez encore, faites-vous plaisir avec la section suivante !

IV-C-2-a - Plus de précisions

Chacun de ces deux processeurs contient une interface qui code et décode les instructions et qui les lance. Derrière l'interface, plusieurs unités exécutent les instructions. Les calculateurs fonctionnent deux fois plus vite que l'interface !

Ces calculateurs sont huit unités de calcul (SP) et deux unités superfonctionnelles (SFU).

À chaque cycle, l'interface choisit un warp prêt à être exécuté.

Pour exécuter toutes les instructions des 32 threads, il faudra quatre cycles. Cependant, vu de l'interface, cela prendra deux cycles.

Pour éviter que l'interface reste inactive pendant un cycle, l'idéal est d'alterner les types de warps : un premier pour les SP, un second pour les SFU.

IV-C-2-b - Limites

Un SM étant composé de huit SP, on sera donc limité à l'exécution de huit blocs en simultané. De plus, l'exécution est limitée à 65 536 blocs et 512 threads par bloc au total.

Vous n'avez pas encore eu un aperçu du temps consacré au calcul en fonction des différents paramètres.

En faisant varier le nombre de blocs de calcul sur un même problème, voici les résultats que l'on peut obtenir, avec de simples opérations d'entrée/sortie dans une table. Un bloc correspond à un thread sur un CPU, que l'on peut affecter à un cœur.

Le processeur utilisé ici est un simple cœur, ses performances en fonction du nombre de threads restent donc stables. S'il s'agissait d'un quad-core, le minimum serait situé à quatre threads.

La carte graphique, une GeForce 8800 GTX, possède seize processeurs, qui ne donnent leur pleine puissance qu'à deux blocs chacun. NVIDIA recommande toutefois d'utiliser au moins une centaine de blocs, afin de pouvoir utiliser la puissance de chipsets plus récents et à venir.

IV-D - CPU

Il exécute uniquement les instructions dans l'ordre assigné, sans parallélisation (sauf architectures multicores et multiCPU, qui nécessitent quand même une action à la conception).

Les instructions sont aussi écrites en mémoire, pour exécution. Cependant, les données avec lesquelles il faudra travailler sont souvent dans la même mémoire !

D'habitude, il ne travaille pas directement en mémoire : les données sont copiées dans des registres puis manipulées et enfin stockées en mémoire.

À l'origine, CPU et mémoire partageaient les mêmes fréquences. Mais le premier a accéléré et la seconde ne l'a pas rattrapé : au point que, si les processeurs actuels lisaient directement dans la mémoire, ils ne seraient utilisés que 10 % du temps.

IV-D-1 - Mémoire cache

C'est pour cela que des caches ont été installés : il s'agit de petites quantités de mémoire, mais très rapide, qui se placent entre le CPU et la mémoire centrale. Ils ne sont utilisés que pour les instructions fréquemment utilisées et les données. Il en existe deux niveaux : L1 et L2, exceptionnellement un troisième, L3, sur les processeurs les plus chers (réservés généralement aux serveurs).

Cependant, ces mémoires très rapides ne sont pas présentes en grandes quantités sur nos CPU, vu leur prix : en moyenne, le mégaoctet de cache coûte 100 fois plus cher que le mégaoctet de RAM ! Le cache fonctionne aussi dix fois plus vite que la RAM, avec un temps d'accès de cinq à dix fois inférieur.

Les caches sont utilisés de manière transparente par le matériel. Ils se font les miroirs des données en mémoire. Ils transportent les données où elles sont nécessaires quand cela est demandé. Ces données ne sont remplacées que quand des données plus urgentes arrivent.

Si les données demandées par le CPU sont disponibles sur le cache, celui-ci les lui envoie, le CPU ne doit pas attendre. Par contre, si elles ne le sont pas, la demande est effectuée en aval, sur des mémoires plus lentes et le CPU doit attendre.

IV-D-2 - Pipelines d'instructions

Supposons qu'un CPU prenne trois cycles pour une multiplication de paires. Combien de temps prendra-t-il pour multiplier n paires ? Nous pourrions dire 3 n cycles. Il est possible de réduire ce nombre.

La multiplication aura lieu dans une ligne de production. Nous pouvons avoir plus d'une paire de nombres en calcul en même temps. Dans ce cas, les multiplications prendront n + 2 cycles.

Notre but, pour atteindre cette vitesse, est de garder le pipeline rempli.

Dans une architecture avec pipelines, il est préférable d'avoir le moins possible de branches.

Moins de branches

do i=start,end
a(i) = b * c(i)
end do

Moins de multiplications

do i=start,end
if( c(i) == 0 )
a(i) = 0
else if( c(i) == 1 )
a(i) = b
else
a(i) = b * c(i)
end if
end do

IV-D-3 - Exécution superscalaire

Les CPU modernes ont plusieurs unités de calcul, qui peuvent effectuer un nombre limité d'instructions en parallèle.

Le matériel examine les instructions pour repérer des opportunités d'optimisation.

Le pipeline

i = i + 1;
j = j + 1;
a = b * c;

Les branches limitent ces opportunités et les unités d'exécution sont laissées en attente pendant l'évaluation des conditions.

Les CPU essayent toutes sortes d'autres astuces, comme la prédiction de branches, l'exécution spéculative ou autres, dont le compilateur et le CPU s'occupent.

V - Les mains dans le cambouis

Il n'y a pas de langage informatique dans lequel vous ne puissiez écrire de mauvais programme.

Si vous ne savez pas ce que votre programme est censé faire, vous feriez bien de ne pas commencer à l'écrire.

(Extraits de Les lois de Murphy).

V-A - Les kernels

Très simplement, un kernel est une fonction exécutée sur le GPU.

Il en existe différents types, qualifiés de :

1. __global__ : exécuté sur le GPU, mais appelé par le CPU ;

2. __device__ : exécuté et appelé par le GPU ;

3. __host__ : exécuté et appelé par le CPU.

Ce dernier qualificatif n'est pas obligatoire : c'est le mode de fonctionnement par défaut.

Un kernel ne s'appelle pas de la même manière qu'une fonction. Voici un appel de fonction :

Appel de fonction

fonction(parametre, parametre);

Mais avant de vous parler de l'appel d'un kernel, il faut que vous compreniez bien le mode de fonctionnement d'un GPU.

Une grille représente la totalité de la tâche à effectuer. Chaque grille peut être divisée en un ou plusieurs blocs, chacun exécutant plusieurs threads.

Un thread sur un GPU n'a pas le même sens qu'un thread sur le CPU.

Sur un GPU, il s'agit de la plus petite subdivision de la tâche à effectuer.

Un appel de kernel se fait en spécifiant deux paramètres entre triples chevrons précédant les paramètres passés au kernel.

kernel <<< nBlocs, threadsParBloc >>> (arguments);

nBlocs est le nombre de subdivisions appliquées à la grille à calculer et est de type dim3 (le cast à partir d'un entier N initialise le dim3 à {N, 1, 1}).
threadsParBloc indique le nombre de threads à exécuter simultanément pour chaque bloc. Ici encore, cette valeur est de type dim3.

Les valeurs à appliquer dépendent simultanément du problème à résoudre (choix des dimensions des blocs) et du matériel utilisé (nombre de threads par bloc). Choisir un nombre de threads supérieur à la quantité nativement supportée entraînera une perte de performances. Cette notation permet ainsi d'adapter dynamiquement le programme aux matériels passés, présents et futurs.

Chaque kernel dispose de variables implicites en lecture seule (toutes de type dim3).

1. blockIdx : index du bloc dans la grille ;

2. threadIdx : index du thread dans le bloc ;

3. blockDim : nombre de threads par bloc (valeur de threadsParBloc du paramétrage du kernel).

La grille est ici considérée comme un seul et unique bloc à une seule dimension.

__global__ void vecAdd(float * A, float * B, float * C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}

int main()
{
// utilisation du kernel
vecAdd<<<1, N>>>(A, B, C);
//     |-> vecteurs additionnés une seule fois
//        |-> nombre de composantes des vecteurs
}

Dans le cas où la grille est sous-divisée en N blocs (tous d'une seule dimension), l'index pourrait être trouvé de la manière suivante.

__global__ void vecAdd(float * A, float * B, float * C)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
C[i] = A[i] + B[i];
}

int main()
{
// utilisation du kernel
const int nThreadsPerBlocks  = 4;
const int nBlocks            = (arraySize / nThreadsPerBlocks) + ( (arraySize % nThreadsPerBlocks) == 0 ? 0 : 1);
vecAdd<<<nBlocks, nThreadsPerBlocks>>>(A, B, C);
}

Les variables doivent être qualifiées, pour définir leur lieu de résidence : voyez la section qui y est réservée.

Les paramètres entre chevrons sont requis, car le kernel est de type __global__. Si le kernel était d'un autre type, les paramètres n'auraient pas dû être précisés !

V-B - Qualifieurs de kernels

V-B-1 - __global__

- Exécuté sur le périphérique.

- Appelable de l'hôte.

- Pas de récursion possible.

- Pas de variables statiques.

- Pas de liste de paramètres variables.

- On ne peut demander leur adresse mémoire.

- Incompatible avec __device__.

- Ne peut rien retourner.

- À l'exécution, on doit préciser la configuration.

- Appel asynchrone (le kernel retourne avant d'avoir effectué les calculs).

- Paramètres stockés dans la mémoire partagée, limités à 256 octets.

- Dure aussi longtemps que le kernel.

V-B-2 - __device__

- Exécuté sur le périphérique.

- Appelable du périphérique.

- Pas de récursion possible.

- Pas de variables statiques.

- Pas de liste de paramètres variables.

- On ne peut demander leur adresse mémoire.

- Incompatible avec __global__.

- Dure aussi longtemps que l'application.

V-B-3 - __host__

- Exécuté sur l'hôte.

- Appelable de l'hôte.

- Appliqué par défaut.

- Compatible avec __device__ (dans ce cas, le kernel pourra être exécuté sur l'hôte et sur le périphérique).

- Incompatible avec __global__.

- Dure aussi longtemps que le kernel.

V-C - Configuration de l'exécution

Ceci n'est requis que pour les kernels __global__ ! Requis signifie bien que l'on ne peut s'en passer, sans quoi rien ne fonctionne (avec, à la clé, beaux plantages) !

Cette configuration doit être passée entre triples chevrons avant les paramètres.

//Définition du kernel
__global__ void func(float * parameter);
//Utilisation du kernel
func <<< Dg, Db, Ns, S >>> (parameter);

V-C-1 - Dg

- Type : dim3.

- Utilité : spécifier la taille et la dimension de la grille (le produit des trois composantes est le nombre de blocs lancés).

- Remarque : z n'est pas encore utilisé et doit valoir 1.

V-C-2 - Db

- Type : dim3.

- Utilité : spécifier la taille et la dimension de chaque bloc (le produit des trois composantes est le nombre de threads par bloc).

V-C-3 - Ns

- Type : size_t.

- Utilité : spécifier le nombre d'octets en mémoire partagée alloués dynamiquement par bloc en plus de la mémoire allouée statiquement.

- Remarque : paramètre optionnel, valeur par défaut : 0.

V-C-4 - S

- Type : cudaStream_t.

- Utilité : spécifier le flux associé.

- Remarque : paramètre optionnel, valeur par défaut : 0.

- La notion de flux sera abordée plus tard : sachez simplement qu'il s'agit d'une suite d'éléments de même type (comme une texture).

V-D - Qualificateurs de variables

V-D-1 - __device__

Cette variable est et restera sur le périphérique. Elle ne vivra pas plus longtemps que l'application et est accessible à tous les threads de la grille et à l'hôte grâce au runtime.

Ce type peut se marier avec un des deux suivants.

V-D-2 - __constant__

Ce type peut être utilisé avec __device__.

La variable restera en mémoire constante. Elle ne vivra pas plus longtemps que l'application et est accessible à tous les threads de la grille et à l'hôte par le runtime.

Ces variables ne peuvent être déclarées que de l'hôte, pas du périphérique !

V-D-3 - __shared__

Ce type peut être utilisé avec __device__.

La variable résidera dans la mémoire partagée et ne survivra pas au bloc. Elle ne sera accessible qu'aux threads du bloc.

Avant que les modifications soient écrites dans la variable et visibles pour tous les autres threads, il faut appeler __syncthreads();. À noter que cet appel ne sert qu'à le garantir, il est possible que les modifications soient visibles avant.

Tableau externe

extern __shared__ float shared[];

Quand la variable est déclarée en tant que tableau externe, comme précédemment, sa taille sera fixée à l'exécution. Toutes les variables déclarées de cette manière ne sont pas contiguës : le premier bit de la première correspond au premier bit des autres, contrairement aux autres langages comme le C ou le C + +.

C'est pourquoi il faut préciser l'offset de début. Pour avoir l'équivalent de ce premier code, il faut écrire le contenu du second.

Code C++

short array0[128];
float array1[64];
int array2[256];

Équivalent CUDA

extern __shared__ char array[];
__device__ void func()        // kernel __device__ ou bien __global__
{
short* array0 = (short*) array;
float* array1 = (float*)&array0[128];
int  * array2 = (int*)  &array1[64];
}

Ceci est le seul moyen d'utiliser le mot-clé extern sur des variables : tous les autres emplois sont interdits.

Ces variables ne peuvent pas être initialisées en même temps que leur déclaration !

Si nous avions utilisé le premier code dans CUDA, en écrivant une valeur dans le premier tableau, une partie de cette variable aurait été imputée au deuxième et au troisième tableau. Ce qui pourrait donner des résultats très aberrants.

V-D-4 - Généralités

Ces paramètres ne sont pas permis sur des unions ou des structures.

En définissant une variable __shared__ ou __constant__, elle sera définie statique.

V-E - Compilation

NVIDIA, dans son immense bonté, nous fournit un compilateur prévu pour CUDA. Celui-ci dispose d'une interface en ligne de commande simple et comparable à celles que nous connaissons, cl, de Visual Studio, ou gcc, l'interface de GCC. Ce compilateur, nvcc, s'occupe de toutes les étapes de la compilation.

Pour pouvoir définir les portions de code spécifiques à ce compilateur, il définit la macro __CUDACC__.

Comme dit précédemment, il s'occupe de toutes les phases de la compilation : l'assemblage, la compilation, et l'édition des liens. Vous pouvez choisir ces parties grâce à la ligne de commande.

Ce compilateur fonctionne très bien avec les Makefiles, c'est d'ailleurs cette technique qui va être ici développée, compatible avec les chaînes de compilation GNU (make) et Microsoft (nmake).

# Précise le compilateur précis à utiliser
ifdef ON_WINDOWS
export compiler-bindir := "a:/program files/microsoft visual studio 9.0/vc/bin"
endif

export NVCC := a:/cuda/bin/nvcc.exe

cpp.obj : cpp.cpp
$(NVCC) -c cpp.cpp $(CFLAGS) -o cpp.obj

c.o : c.c
$(NVCC) -c c.c $(CFLAGS) -o c.obj

cu.o : cu.cu
$(NVCC) -c cu.cu $(CFLAGS) -o cu.obj

OBJECTS = cpp.obj c.obj cu.obj

all : $(OBJECTS)
$(NVCC) $(OBJECTS) $(LDFLAGS) -o app.exe

clean :
$(RM) $(OBJECTS)

Ce Makefile doit être utilisé après avoir appelé le script vsvars.bat s'il est utilisé avec Visual Studio !

Si vous utilisez un make d'origine GNU, vous pouvez utiliser ce Makefile

# Précise le compilateur précis à utiliser
ifdef ON_WINDOWS
export compiler-bindir := "a:/program files/microsoft visual studio 9.0/vc/bin"
endif

export NVCC := a:/cuda/bin/nvcc.exe

%.o : %.cpp
$(NVCC) -c %^ $(CFLAGS) -o $@
$(NVCC) -M %^ $(CFLAGS) > $@.dep

%.o : %.c
$(NVCC) -c %^ $(CFLAGS) -o $@
$(NVCC) -M %^ $(CFLAGS) > $@.dep

%.o : %.cu
$(NVCC) -c %^ $(CFLAGS) -o $@
$(NVCC) -M %^ $(CFLAGS) > $@.dep

include $(wildcard *.dep) /dev/null

all : $(OBJECTS)
$(NVCC) $(OBJECTS) $(LDFLAGS) -o app.exe

clean :
$(RM) $(OBJECTS) *.dep

Vous pouvez aussi utiliser la ligne de commande directement :

nvcc -c cu.cu -o cu.obj
nvcc cu.obj -o app.exe

Vous pouvez aussi décider que le code CUDA sera exécuté sur le processeur, qui émulera alors un GPU. Il suffit d'ajouter emu=1 à la ligne de commande, comme ceci :

make emu=1

VI - Le modèle de programmation

Pour un développeur CUDA, l'ordinateur consiste en un ou plusieurs hôtes, un traditionnel CPU et un ou plusieurs périphériques, des non moins traditionnels GPU, des processeurs massivement parallèles.

Dans les applications modernes, certaines parties utilisent du calcul qui peut facilement devenir parallèle, sans aucun problème. Ces parties peuvent être déportées de l'hôte vers le périphérique.

VI-A - Parallélisme des données

Les applications actuelles qui doivent traiter de grandes quantités de données prennent beaucoup de temps à l'exécution. Ce temps pourrait être réduit en parallélisant les opérations : des phénomènes physiques peuvent être calculés indépendamment les uns des autres, des images à analyser peuvent être découpées en portions et un flux vidéo peut être découpé image par image.

La parallélisation des données réfère à la propriété du programme de gérer parallèlement et indépendamment ces instructions arithmétiques.

Par exemple, pour des multiplications de matrices de taille 1000 x 1000, il s'agit de 1 000 000 de multiplications, sans rapport les unes avec les autres, qui peuvent donc être parallélisées sans problème. Un GPU peut fortement améliorer les performances en exécutant toutes ces opérations simultanément.

VI-B - Structure du programme

Un programme CUDA est constitué d'une partie qui s'exécute sur l'hôte et d'une partie qui s'exécute sur le périphérique.

Les phases peu ou pas parallèles sont exécutées sur l'hôte.

Les phases massivement parallèles sont exécutées sur le périphérique.

Le programme peut tenir en un seul fichier, comprenant ces deux phases et environnements. Le compilateur se charge de les séparer : le code pour l'hôte est du standard C ANSI/ISO et est compilé par le compilateur principal du système, il sera lancé comme un simple processus. Le code pour le périphérique est aussi écrit en C ANSI/ISO, avec quelques extensions CUDA, mais il est compilé par NVCC et sera exécuté sur le périphérique.

Les kernels génèrent généralement beaucoup de threads pour exploiter au mieux le parallélisme des données.

Dans notre exemple de produit matriciel, il y a autant de threads que de cellules dans la matrice résultante. Chacun de ces threads prend, généralement, très peu de cycles, vu le peu de tâches qui leur sont demandées.

L'exécution commence avec le CPU, qui prépare l'appel au kernel. Le GPU prend le relais pour le kernel, qui sera, lui, massivement multithread. Quand le kernel a fini sa tâche, il renvoie le résultat au CPU et son exécution continue.

VI-C - L'exemple : la multiplication de matrices carrées

Pour commencer par clarifier la situation, voici le fonctionnement que décrira notre programme :

1. CPU : initialisation des matrices M, N et P, toutes carrées ;

2. CPU : remplissage des matrices d'entrée M et N ;

3. GPU : calcul du produit matriciel de M et de N, dont le résultat est stocké sur P ;

4. CPU : écriture de la matrice P ;

5. CPU : nettoyage de la mémoire et fin de l'exécution du programme.

Nous avons vu les différents types de mémoire, mais pas la manière d'y accéder. Or, cela sera nécessaire pour permettre l'exécution du programme. Nous allons ici nous concentrer sur l'utilisation de la mémoire globale, le but étant de montrer le fonctionnement d'un programme CUDA et non d'optimiser au maximum une application.

Ces fonctions se trouvent, heureusement, dans l'API CUDA. Leur nom est très recherché : cudaMalloc() et cudaFree().

Voici un bref exemple d'utilisation de ces deux fonctions, très proches des fonctions malloc() et free() du C. On considère que Width est le nombre de lignes et de colonnes de la matrice pour laquelle on crée l'espace mémoire.

float *Md;
float *Nd;
float *Pd;
const int size = Width * Width * sizeof(float);

cudaMalloc( (void**) & Md, size);
cudaMalloc( (void**) & Nd, size);
cudaMalloc( (void**) & Pd, size);

cudaFree  (            Md      );
cudaFree  (            Nd      );
cudaFree  (            Pd      );

cudaMalloc() prend deux paramètres, pour définir la mémoire à allouer en mémoire globale :

1. L'adresse d'un pointeur vers la mémoire allouée ;

2. La taille de la mémoire à allouer.

cudaFree() ne prend qu'un paramètre, pour désallouer cette mémoire en mémoire globale:

1. Un pointeur vers la mémoire à désallouer.

Une fois que le programme a alloué sa mémoire, il peut demander les données des matrices à stocker en mémoire.

Ceci s'obtient avec une fonction de copie de mémoire : cudaMemcpy(). Cette fonction requiert quatre paramètres :

1. Un pointeur vers les données source à copier ;

2. La destination des données ;

3. Le nombre d'octets à copier ;

4. Le type de mémoire vers laquelle copier.

Concernant le quatrième paramètre, il peut prendre une de ces valeurs :

- cudaMemcpyHostToDevice : copie de l'hôte vers le périphérique ;

- cudaMemcpyHostToHost : copie de l'hôte vers l'hôte ;

- cudaMemcpyDeviceToHost : copie du périphérique vers l'hôte ;

- cudaMemcpyDeviceToDevice : copie du périphérique vers le périphérique.

Voici les appels réalisés pour copier les matrices sur lesquelles nous allons travailler et pour envoyer le résultat à l'endroit souhaité. M, N, P, Md, Nd, Pd et size gardent leur valeur précédente.

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);

cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

Tous ces transferts de mémoire sont asynchrones !

Maintenant que nous savons ce que nous pouvons faire avec la mémoire et comment le faire, nous pouvons commencer l'implémentation de notre exemple. Normalement, à ce stade de votre étude de CUDA, vous devriez pouvoir écrire correctement ce kernel (ce qui est fortement recommandé, il s'agit d'un exercice comme un autre), et voici la correction.

__global__ void MatrixMulKernel(float * Md, float * Nd, float * Pd, int Width)
{
// identifiant de thread à deux dimensions, comme la matrice
int tx = threadIdx.x;
int ty = threadIdx.y;
// Pvaleur sert au stockage de la valeur calculée par le thread
float Pvaleur = 0;
for (int i = 0; i < Width; ++i)
{
float MdElement = Md[ty * Width + i];
float NdElement = Nd[i  * Width + tx];
Pvaleur        += MdElement * NdElement;
}
// écrit la valeur calculée dans la matrice de résultat
// chaque thread ne peut écrire qu'une valeur !
Pd[ty * Width + tx] = Pvaleur;
}

Et voici l'utilisation de ce kernel, que vous devriez aussi pouvoir écrire.

void MatrixMulOnDevice(float * M, float * N, float * P, int Width)
{
//calcul de la taille des matrices
int size = Width * Width * sizeof(float);

//allocation des matrices et leur remplissage
cudaMalloc(Md, size);
cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice) ;
cudaMalloc(Nd, size);
cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);

//allocation de la matrice de résultat
cudaMalloc(Pd, size);

//multiplication d'une seule matrice
dim3 dimGrid(1, 1);
//matrice carrée
dim3 dimBlock(Width, Width);

//produit matriciel proprement dit
MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);

//récupération du résultat du calcul
cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

//destruction des matrices, désormais inutilisées
cudaFree(Md);
cudaFree(Nd);
cudaFree(Pd);
}

Pour que ce code compile, vous devez inclure les fichiers cuda.h et cuda_runtime.h.

VI-D - Séparer les opérations à effectuer

Comme souvent dit plus haut, le périphérique exécute beaucoup de calculs en même temps. Il est donc bien nécessaire de découper ses calculs en autant de petits morceaux, algorithmiquement identiques.

Prenons un exemple : le traitement de données sismiques. Il s'agit d'analyser une image, à trois dimensions, pixel par pixel, pour vérifier l'évolution de la situation. Nous allons nous concentrer sur l'imagerie de Kirchhoff.

Dans ce cas, une grille représente une image à analyser, avec deux dimensions, qui correspondent aux largeur et hauteur de l'image. Un bloc aura donc l'abscisse et l'ordonnée déjà fixées, la seule dimension à encore faire varier est la cote : voici donc le nombre de threads par bloc. Chaque thread aura donc ses trois coordonnées définies par ses places dans la grille et dans le bloc, il comparera ce pixel avec celui de l'image précédente.

Ceci correspond, plus ou moins, à l'algorithme utilisé sur un CPU : il est constitué de trois boucles imbriquées, chacune faisant varier une coordonnée. Basiquement, ces boucles ne sont pas parallélisées, mais cela peut être effectué sans problème, ce qui améliore sensiblement les performances.

VII - Conclusions

Nous avons fini une très brève introduction aux possibilités offertes par CUDA. Nous n'avons vu que le strict nécessaire pour commencer à écrire des programmes CUDA et les compiler.

Vous avez pu vous rendre compte de la simplicité d'écrire du code avec CUDA. Cette approche, relativement haut niveau, ne permet pas au premier abord une optimisation approfondie. Pour ce genre d'exercice, il faut se tourner vers une solution plus proche du matériel, comme l'antique CTM. Cependant, optimiser à ce niveau augmente considérablement le temps de développement (si vous avez déjà utilisé l'assembleur, vous savez ce que c'est).

VII-A - Et chez AMD/ATI ?

Cette société a été la première à dégainer avec CTM, mais elle a pris du retard : son langage haut-niveau a été lancé en 2008, soit deux ans après CUDA. Donc, le langage Brook a été repris. Mais il produisait du code OpenGL. AMD l'a donc amélioré pour qu'il produise du CAL (proche de l'assembleur). Et s'est arrêté là.

Dernièrement, ils ont abandonné Brook+ (leur version de Brook) pour OpenCL, qui est en passe de devenir un standard accepté par tous (y compris NVIDIA, dont le support est disponible en bêta privée pour le moment).

Il n'est pas possible d'utiliser le GPU comme processeur sans le remarquer à toutes les lignes : il faut commencer par une initialisation, qui doit préciser la version des Shaders à utiliser. Certaines commandes existent en version DirectX9 et DirectX10.

La documentation n'est pas là pour aider : aussi brouillonne que l'interface, elle est plus qu'illisible. Principal grief : l'emploi des noms pour les GPU. Nous entendons parler de Radeon HD 2900, de R600, de Pele. Qui ne sont, en fait, que les mêmes GPU…

De quoi rebuter facilement du monde. Même si l'architecture peut être plus performante et permet une double précision (FP64) depuis plus longtemps.

Dernier point : le SDK pour AMD Stream est réservé à ses GPU, réservés aux professionnels. La firme se détourne complètement des GPU grand public, contrairement à son opposant.

VII-B - Intégration à Visual Studio

CUDA est prévu pour le compilateur de Visual Studio, sous Windows, alors que NVIDIA ne propose strictement aucun moyen d'intégrer le SDK dans l'IDE. C'est pourquoi Kaiyong Zhao a créé un modèle de projet pour cet IDE. Ce modèle est compatible avec Visual Studio 2005 (8.0) et 2008 (9.0).

Il vous suffit d'aller sur le site web du projet pour télécharger le modèle, puis de l'installer. Ainsi, un nouveau type de projet sera disponible à la création.

Cependant, cela ne vous apportera pas la coloration syntaxique du code. Pour ce faire, ouvrez le fichier C:\Program Files\Microsoft Visual Studio 9.0\Common7\IDE\usertype.dat (chemin à modifier selon votre installation). S'il n'existe pas, créez-le. Ajoutez-y le contenu de C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\doc\ syntax_highlighting\visual_studio_8\usertype.dat. Il s'agit de l'ensemble des mots-clés, des types, des variables prédéfinies et des fonctions mathématiques de base de CUDA. Ce fichier convient aux versions 7.0, 7.1, 8.0 et 9.0.

Maintenant, il faut que Visual Studio utilise une coloration syntaxique, semblable à celle du C et du C + +, avec les quelques mots-clés que nous venons d'ajouter, pour les fichiers CUDA.

Pour les versions 7.0 et 7.1, utilisez le fichier C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\doc\ syntax_highlighting\visual_studio_7\install_cuda_highlighting_vs7.reg.

Pour les autres versions, Outils > Options > Éditeur de texte > Extension de fichier, ajoutez les extensions cu et cuh, avec l'éditeur Visual C + +.

Et voilà ! Visual Studio est maintenant fin prêt à créer de nouveaux projets CUDA et à les mettre en couleur comme il le faut !

Il existe une autre solution : le fichier de règles pour CUDA. Il s'utilise comme tous les autres fichiers de règles. Placez ce fichier dans le dossier Microsoft Visual Studio 9.0\VC\VCProjectDefaults\ : cuda.rules. De cette manière, tous les projets disposeront des règles pour construire des fichiers .cu.

VII-C - Avec d'autres environnements de compilation

CUDA étant compatible avec une pléthore de compilateurs, pourquoi serait-il impossible de l'utiliser avec d'autres chaînes de compilation ? Ici ne sera montré qu'un bref exemple avec qmake.

qmake est la chaîne de compilation utilisée avec la bibliothèque C + + Qt, éditée par Nokia. Voici un petit script qui vous permettra d'utiliser le framework Qt et la bibliothèque C for CUDA sans problème de génération du côté Qt. Il vous suffit d'ajouter à la variable CUDA_SOURCES les fichiers qui doivent passer par nvcc.

Avec qmake

win32 {
INCLUDEPATH += $(CUDA_INC_DIR)
QMAKE_LIBDIR += $(CUDA_LIB_DIR)
LIBS += -lcudart

cuda.output = $$OBJECTS_DIR/${QMAKE_FILE_BASE}_cuda.obj
cuda.commands  = $(CUDA_BIN_DIR)/nvcc.exe -c -Xcompiler $$join(QMAKE_CXXFLAGS,",") $$join(INCLUDEPATH,'" -I "','-I "','"')
cuda.commands += ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT}
}
unix {
# auto-detect CUDA path
CUDA_DIR = $$system(which nvcc | sed 's,/bin/nvcc$,,')
INCLUDEPATH += $$CUDA_DIR/include
QMAKE_LIBDIR += $$CUDA_DIR/lib
LIBS += -lcudart

cuda.output = ${OBJECTS_DIR}${QMAKE_FILE_BASE}_cuda.obj
cuda.commands  = nvcc -c -Xcompiler $$join(QMAKE_CXXFLAGS,",") $$join(INCLUDEPATH,'" -I "','-I "','"')
cuda.commands += ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT}
cuda.depends  = nvcc -M -Xcompiler $$join(QMAKE_CXXFLAGS,",") $$join(INCLUDEPATH,'" -I "','-I "','"')
cuda.depends += ${QMAKE_FILE_NAME} | sed "s,^.*: ,," | sed "s,^ *,," | tr -d '\\\n'
}
cuda.input = CUDA_SOURCES
QMAKE_EXTRA_UNIX_COMPILERS += cuda

VII-D - Déploiement

Vous ne devrez déployer strictement aucune DLL ou autre à côté de votre application : CUDA est intégré aux pilotes graphiques, depuis CUDA 1.1 et les pilotes de génération 177 pour GeForce. Par contre, pour que votre application fonctionne, il faut que le client dispose d'une carte graphique compatible et de pilotes supportant CUDA ; ou alors que vous compiliez votre application en mode émulation.

Cependant, ceci n'est valable que si vous n'utilisez que le driver de CUDA. Si vous utilisez le runtime (présenté ici), vous devrez aussi en distribuer la DLL (cudart.dll).

Si le client ne dispose pas d'une carte graphique compatible, vous ne devrez pas livrer de DLL particulière si vous vous limitez à CUDA : si vous utilisez une autre bibliothèque, vous devrez aussi en livrer la DLL d'émulation (suffixée par _emu).

VII-E - Quels fichiers inclure ?

Ici, je n'ai parlé que du runtime : il est composé de deux fichiers, cuda_runtime.h et cuda_runtime_api.h. Tous deux sont nécessaires pour pouvoir utiliser l'ensemble des fonctionnalités du runtime.

VIII - Divers

VIII-A - Références

- ECE Illinois - cours 498 AL : Programming Massively Parallel Processors

- Tesla 10 & Cuda 2.0

- CUDA, Supercomputing for the Masses

- The End of the CPU ?

- GPU Workshop

- CUDA Textbook, disponible sur le site du MIT

- NVIDIA CUDA Programming Guide (versions 2.1 et 2.2), distribué avec le SDK

- NVIDIA NVCC Guide (versions 2.1 et 2.2), distribué avec le SDK

VIII-B - Voir aussi

VIII-B-1 - GPGPU

- La FAQ GPGPU

- BrookGPU (Stanford)

- Close to Metal (ATI)

- Stream (ATI/AMD)

- CUDA (NVIDIA)

- OpenCL (Apple/Khronos)

- Larrabee (Intel)

VIII-B-2 - CUDA

- Site officiel

- Téléchargements

- Produits compatibles

- CUDA Visual Studio Wizard

VIII-B-3 - Bibliothèques

- BLAS

- Transformée rapide de Fourier (FFT)

VIII-B-4 - Matériel

- DRAM

- DRAM

VIII-B-5 - Autres

- Précision des nombres à virgule flottante (FP32 et FP64)

- Matrices

- Produit matriciel

VIII-C - Remerciements

Un tout grand merci à raptor70, à gorgonite, à ange_blond, à TanEk et surtout à IrmatDen, Matthieu Brucher, et à 3DArchi pour sa courageuse relecture, et Étienne Bougoin, ainsi que Sébastien Flochlay, pour leurs commentaires, encouragements et précisions, sans lesquels l'article ne serait pas celui qu'il est maintenant !

Merci aussi à tous les lecteurs qui ont relevé les quelques fautes qui se trouvaient encore dans l'article : L0ur5.