- un ensemble d'interfaces logicielles de bas niveau ( API) pour créer des jeux et autres applications multimédias performantes. Comprend un support pour des performances élevées 2D- et 3D-périphériques graphiques, audio et d'entrée.

Direct3D (D3D) - interface pour la sortie en trois dimensions primitives(corps géométriques). Inclus dans .

Opengl(de l'anglais. Ouvrir la bibliothèque graphique, littéralement - une bibliothèque graphique ouverte) est une spécification qui définit une interface de programmation multiplateforme indépendante d'un langage de programmation pour écrire des applications à l'aide d'infographies bidimensionnelles et tridimensionnelles. Comprend plus de 250 fonctions pour dessiner des scènes 3D complexes à partir de simples primitives. Il est utilisé dans la création de jeux vidéo, la réalité virtuelle, la visualisation dans la recherche scientifique. Sur la plateforme les fenêtres rivalise avec .

OpenCL(de l'anglais. Langage informatique ouvert, littéralement - un langage informatique ouvert) - cadre(framework d'un système logiciel) pour l'écriture de programmes informatiques liés à des calculs parallèles sur divers graphiques ( GPU) et ( ). Dans le cadre OpenCL comprend un langage de programmation et une interface de programmation d'applications ( API). OpenCL fournit un parallélisme au niveau des instructions et des données et est une implémentation de la technique GPGPU.

GPGPU(abrégé de l'anglais. Unités de traitement de graphiques à usage général, au sens propre - GPU usage général) est une technique d'utilisation d'un processeur graphique-carte vidéo pour les calculs généraux qui sont généralement effectués.

Ombrage(eng. ombrage) - un programme d'ombrage sur des images de synthèse, utilisé dans les graphiques en trois dimensions pour déterminer les paramètres finaux d'un objet ou d'une image. Comprend généralement une complexité arbitraire décrivant l'absorption et la diffusion de la lumière, le mappage de texture, la réflexion et la réfraction, l'ombrage, le déplacement de surface et les effets de post-traitement. Les surfaces complexes peuvent être rendues à l'aide de formes géométriques simples.

Le rendu(eng. le rendu) - visualisation, en infographie, du processus d'obtention d'une image à partir d'un modèle à l'aide d'un logiciel.

SDK(abrégé de l'anglais. Kit de développement logiciel) - un ensemble d'outils de développement logiciel.

CPU(abrégé de l'anglais. Unité centrale de traitement, littéralement - dispositif informatique central / principal / principal) - central (micro); un dispositif qui exécute les instructions de la machine ; un élément matériel chargé d'effectuer des opérations de calcul (spécifiées par le système d'exploitation et le logiciel d'application) et de coordonner le travail de tous les appareils.

GPU(abrégé de l'anglais. Unité de traitement graphique, littéralement - un périphérique informatique graphique) - un processeur graphique ; Un appareil ou une console de jeu distinct qui effectue le rendu graphique (rendu). Les GPU modernes gèrent et restituent très efficacement les graphiques informatiques. Le processeur graphique des adaptateurs vidéo modernes est utilisé comme accélérateur pour les graphiques en trois dimensions, mais il peut également être utilisé dans certains cas pour des calculs ( GPGPU).

Problèmes CPU

Pendant longtemps, une augmentation des performances des traditionnels était principalement due à une augmentation séquentielle de la fréquence d'horloge (environ 80% des performances étaient déterminées par la fréquence d'horloge) avec une augmentation simultanée du nombre de transistors sur un cristal . Cependant, une nouvelle augmentation de la fréquence d'horloge (à une fréquence d'horloge de plus de 3,8 GHz, les puces surchauffent tout simplement !) se heurte à un certain nombre de barrières physiques fondamentales (puisque le processus technologique est très proche de la taille d'un atome : et la taille d'un atome de silicium est d'environ 0,543 nm) :

Premièrement, à mesure que la taille du cristal diminue et que la fréquence d'horloge augmente, le courant de fuite des transistors augmente. Cela entraîne une augmentation de la consommation d'énergie et une augmentation des émissions de chaleur ;

Deuxièmement, les avantages des vitesses d'horloge plus élevées sont partiellement compensés par la latence de la mémoire car les temps d'accès à la mémoire ne correspondent pas aux vitesses d'horloge croissantes ;

Troisièmement, pour certaines applications, les architectures séquentielles traditionnelles deviennent inefficaces avec des fréquences d'horloge croissantes en raison du "goulet d'étranglement de von Neumann" - une limitation des performances résultant d'un flux séquentiel de calculs. Dans le même temps, les délais de transmission des signaux résistifs-capacitifs augmentent, ce qui constitue un goulot d'étranglement supplémentaire associé à une augmentation de la fréquence d'horloge.

Développement GPU

Parallèlement au développement de GPU:

…

novembre 2008 - Intelligence introduit une ligne de 4-core Intel Core i7 basé sur une microarchitecture de nouvelle génération Nehalem... Les processeurs fonctionnent à une fréquence d'horloge de 2,6 à 3,2 GHz. Fabriqué selon la technologie de processus 45 nm.

Décembre 2008 - expéditions d'un 4-core AMD Phenom II 940(nom de code - Deneb). Fonctionne à une fréquence de 3 GHz, fabriqué à l'aide de la technologie de processus 45 nm.

…

Mai 2009 - entreprise DMLA introduit la version GPU ATI Radeon HD 4890 avec une fréquence d'horloge de base augmentée de 850 MHz à 1 GHz. C'est le premier graphique Processeur 1 GHz. La puissance de calcul de la puce, du fait de l'augmentation de la fréquence, est passée de 1,36 à 1,6 téraflops. Le processeur contient 800 (!) cœurs de calcul, prend en charge la mémoire vidéo GDDR5, DirectX 10.1, ATI CrossFireX et toutes les autres technologies inhérentes aux modèles de cartes vidéo modernes. La puce est fabriquée sur la base de la technologie 55 nm.

Les principales différences GPU

Caractéristiques distinctives GPU(comparé à ) sommes:

- architecture, visant au maximum à augmenter la vitesse de calcul des textures et des objets graphiques complexes ;

- puissance de crête typique GPU beaucoup plus élevé que ;

- grâce à l'architecture de convoyage spécialisée, GPU beaucoup plus efficace dans le traitement des informations graphiques que.

"Crise du genre"

"Crise du genre" pour mûri en 2005 - c'est alors qu'ils sont apparus. Mais, malgré le développement de la technologie, l'augmentation de la productivité des diminué de façon marquée. performances en même temps GPU continuer à grandir. Ainsi, en 2003, cette idée révolutionnaire s'est cristallisée - utiliser la puissance de calcul du graphique... Les GPU sont devenus activement utilisés pour le calcul « non graphique » (simulation physique, traitement du signal, mathématiques/géométrie computationnelles, opérations de base de données, biologie computationnelle, économie computationnelle, vision par ordinateur, etc.).

Le problème principal était qu'il n'y avait pas d'interface standard pour la programmation GPU... Développeurs utilisés Opengl ou Direct3D mais c'était très pratique. société NVIDIA(l'un des plus grands fabricants de processeurs graphiques, multimédias et de communication, ainsi que de processeurs multimédias sans fil ; fondé en 1993) a entrepris le développement d'une norme unifiée et pratique - et a présenté la technologie CUDA.

Comment ça a commencé

2006 - NVIDIA démontre CUDA™; le début d'une révolution informatique GPU.

2007 - NVIDIA architecture des versions CUDA(version originale SDK CUDA a été introduit le 15 février 2007); nomination "Meilleure nouveauté" du magazine Science populaire et "Choix des lecteurs" de la publication HPCWire.

2008 - technologie NVIDIA CUDA remporté dans la catégorie « Excellence technique » de Magazine PC.

Quoi CUDA

CUDA(abrégé de l'anglais. Architecture de périphérique unifiée de calcul, littéralement - architecture informatique unifiée des appareils) - architecture (un ensemble de logiciels et de matériel) qui permet GPU calculs à usage général, tandis que GPU agit en fait comme un coprocesseur puissant.

La technologie NVIDIA CUDA™ Est le seul environnement de développement dans le langage de programmation C, qui permet aux développeurs de créer des logiciels pour résoudre des problèmes de calcul complexes en moins de temps, grâce à la puissance de traitement des GPU. Des millions de personnes travaillent déjà dans le monde GPU avec le soutien CUDA, et des milliers de programmeurs utilisent déjà des outils (gratuits !) CUDA pour accélérer les applications et résoudre les tâches les plus complexes et les plus gourmandes en ressources - de l'encodage vidéo et audio à l'exploration pétrolière et gazière, la modélisation de produits, l'imagerie médicale et la recherche scientifique.

CUDA donne au développeur la possibilité, à sa discrétion, d'organiser l'accès au jeu d'instructions de l'accélérateur graphique et de gérer sa mémoire, d'y organiser des calculs parallèles complexes. Accélérateur graphique avec support CUDA devient une puissante architecture ouverte programmable comme aujourd'hui. Tout cela fournit au développeur un accès de bas niveau, distribué et à haut débit aux équipements, ce qui rend CUDA une base nécessaire lors de la création d'outils sérieux de haut niveau tels que des compilateurs, des débogueurs, des bibliothèques mathématiques, des plates-formes logicielles.

Uralsky, grand spécialiste de la technologie NVIDIA comparant GPU et , le dit : " Est un SUV. Il conduit n'importe quand, n'importe où, mais pas très vite. UNE GPU Est une voiture de sport. Sur une mauvaise route, il n'ira tout simplement nulle part, mais donnera une bonne couverture - et il montrera toute sa vitesse, dont le SUV n'a jamais rêvé ! .. ".

Capacités technologiques CUDA

Depuis des décennies, la loi de Moore est en vigueur, selon laquelle le nombre de transistors sur une puce doublera tous les deux ans. Cependant, c'était en 1965, et au cours des 5 dernières années, l'idée du multicœur physique dans les processeurs grand public a commencé à se développer rapidement : en 2005, Intel a introduit le Pentium D et AMD - l'Athlon X2. Ensuite, les applications utilisant 2 noyaux pourraient être comptées sur les doigts d'une main. Cependant, la prochaine génération de processeurs Intel, qui a fait une révolution, avait exactement 2 cœurs physiques. De plus, en janvier 2007, la série Quad est apparue, au même moment où Moore lui-même admettait que sa loi cesserait bientôt de fonctionner.

Et maintenant? Les processeurs double cœur même dans les systèmes de bureau à petit budget, et 4 cœurs physiques sont devenus la norme en seulement 2-3 ans. La fréquence des processeurs n'augmente pas, mais l'architecture s'améliore, le nombre de cœurs physiques et virtuels augmente. Pourtant, l'idée d'utiliser des adaptateurs vidéo équipés de dizaines voire de centaines de "blocs" de calcul existe depuis longtemps.

Et bien que les perspectives de calcul avec les GPU soient énormes, la solution la plus populaire - Nvidia CUDA est gratuite, a beaucoup de documentation et, en général, est très simple à mettre en œuvre, il n'y a pas beaucoup d'applications utilisant cette technologie. Fondamentalement, ce sont toutes sortes de calculs spécialisés, auxquels l'utilisateur ordinaire ne se soucie pas dans la plupart des cas. Mais il existe également des programmes conçus pour l'utilisateur de masse, et nous en parlerons dans cet article.

Pour commencer, un peu sur la technologie elle-même et avec quoi elle est mangée. Parce que Lors de la rédaction d'un article, je suis guidé par un large éventail de lecteurs, puis j'essaierai de l'expliquer dans un langage accessible sans termes complexes et assez brièvement.

CUDA(Anglais Compute Unified Device Architecture) est une architecture logicielle et matérielle qui permet le calcul à l'aide de processeurs graphiques NVIDIA prenant en charge la technologie GPGPU (calcul arbitraire sur cartes vidéo). L'architecture CUDA est apparue pour la première fois sur le marché avec la sortie de la puce NVIDIA G80 de huitième génération et est présente dans toutes les séries de puces graphiques suivantes utilisées dans les familles d'accélérateurs GeForce, Quadro et Tesla. (c) Wikipédia.org

Les flux entrants sont traités indépendamment les uns des autres, c'est-à-dire parallèle.

De plus, il y a une division en 3 niveaux :

Grille- coeur. Contient un tableau de blocs à une / deux / trois dimensions.

Bloquer- contient de nombreux threads (thread). Les flux de blocs différents ne peuvent pas interagir les uns avec les autres. Pourquoi avez-vous eu besoin d'introduire des blocs ? Chaque bloc est essentiellement responsable de sa propre sous-tâche. Par exemple, une grande image (qui est une matrice) peut être divisée en plusieurs parties plus petites (matrices) et fonctionner en parallèle avec chaque partie de l'image.

Fil- flux. Les threads d'un même bloc peuvent interagir soit via la mémoire partagée, qui, soit dit en passant, est beaucoup plus rapide que la mémoire globale, ou via des outils de synchronisation de threads.

Chaîne- c'est l'union des flux en interaction, pour tous les GPU modernes, la taille de Warp est de 32. Vient ensuite demi-chaîne, qui est la moitié d'une warp'a, puisque l'accès à la mémoire se fait généralement séparément pour la première et la seconde moitié d'une chaîne.

Comme vous pouvez le voir, cette architecture est idéale pour la parallélisation des tâches. Et bien que la programmation se fasse en langage C avec quelques restrictions, en réalité ce n'est pas si simple, puisque tout ne peut pas être parallélisé. Il n'y a pas non plus de fonctions standard pour générer des nombres aléatoires (ou d'initialisation) ; tout cela doit être implémenté séparément. Et bien qu'il existe suffisamment d'options toutes faites, tout cela n'apporte pas de joie. La possibilité d'utiliser la récursivité est relativement nouvelle.

Pour plus de clarté, un petit programme de console (pour minimiser le code) a été écrit qui effectue des opérations avec deux tableaux flottants, c'est-à-dire avec des valeurs non entières. Pour les raisons ci-dessus, l'initialisation (remplissage du tableau avec diverses valeurs arbitraires) a été effectuée par le CPU. Ensuite, 25 opérations différentes ont été effectuées avec les éléments correspondants de chaque tableau, les résultats intermédiaires ont été écrits dans le troisième tableau. La taille du tableau a changé, les résultats sont les suivants :

Au total, 4 tests ont été effectués :

1024 éléments dans chaque tableau :

On voit clairement qu'avec un si petit nombre d'éléments, les calculs parallèles ont peu de sens, car les calculs eux-mêmes sont beaucoup plus rapides que leur préparation.

4096 éléments dans chaque tableau :

Et maintenant, vous pouvez voir que la carte vidéo effectue des opérations sur les matrices 3 fois plus rapidement que le processeur. De plus, le temps d'exécution de ce test sur une carte vidéo n'a pas augmenté (une légère diminution du temps peut être attribuée à une erreur).

Il y a maintenant 12288 éléments dans chaque tableau :

L'écart entre la carte vidéo a doublé. Encore une fois, notez que le temps d'exécution sur la carte vidéo a augmenté.
de manière insignifiante, mais sur le processeur plus de 3 fois, c'est-à-dire proportionnel à la complexité de la tâche.

Et le dernier test - 36864 éléments dans chaque tableau :

Dans ce cas, l'accélération atteint des valeurs impressionnantes - presque 22 fois plus rapides sur une carte vidéo. Et encore une fois, le temps d'exécution sur la carte vidéo a augmenté de manière insignifiante et sur le processeur - les 3 fois prescrits, ce qui est encore une fois proportionnel à la complication de la tâche.

Si vous continuez à compliquer les calculs, la carte vidéo gagne de plus en plus. Bien que l'exemple soit quelque peu exagéré, il montre en général clairement la situation. Mais comme mentionné ci-dessus, tout ne peut pas être parallélisé. Par exemple, calculer pi. Il n'y a que des exemples écrits en utilisant la méthode Monte Carlo, mais la précision des calculs est de 7 décimales, c'est-à-dire flotteur régulier. Afin d'augmenter la précision des calculs, une longue arithmétique est nécessaire, mais ici des problèmes surviennent, car c'est très, très difficile à mettre en œuvre efficacement. Sur Internet, je n'ai pas pu trouver d'exemples utilisant CUDA et calculant Pi à 1 million de décimales. Des tentatives ont été faites pour écrire une telle application, mais la méthode la plus simple et la plus efficace pour calculer pi est l'algorithme de Brent-Salamin ou la formule de Gauss. Dans le SuperPI bien connu, très probablement (à en juger par la vitesse de travail et le nombre d'itérations), la formule de Gauss est utilisée. Et à en juger par
le fait que SuperPI soit monothread, le manque d'exemples pour CUDA et l'échec de mes tentatives, il est impossible de paralléliser efficacement le calcul de Pi.

À propos, vous pouvez voir comment, au cours du processus de calcul, la charge sur le GPU augmente, ainsi que l'allocation de mémoire.

Passons maintenant aux avantages plus pratiques de CUDA, à savoir les programmes actuellement existants utilisant cette technologie. Pour la plupart, il s'agit de toutes sortes de convertisseurs et d'éditeurs audio/vidéo.

Nous avons utilisé 3 fichiers vidéo différents pour les tests :

* L'histoire du film Avatar - 1920x1080, MPEG4, h.264.

* Série "Lie to me" - 1280x720, MPEG4, h.264.

* Série "Il fait toujours beau à Philadelphie" - 624x464, xvid.

Le conteneur et la taille des deux premiers fichiers étaient mkv et 1,55 Go, et le dernier était .avi et 272 Mo.

Commençons par un produit très sensationnel et populaire - Badaboum... La version utilisée était - 1.2.1.74 ... Le coût du programme est $29.90 .

L'interface du programme est simple et intuitive - à gauche, nous sélectionnons le fichier ou le disque source et à droite - le périphérique requis pour lequel nous encoderons. Il existe également un mode utilisateur, dans lequel les paramètres sont définis manuellement, et il a été utilisé.

Pour commencer, considérons la rapidité et l'efficacité avec lesquelles la vidéo est encodée "en elle-même", c'est-à-dire à la même résolution et à peu près la même taille. La vitesse sera mesurée en fps, et non en temps écoulé - il est plus pratique à la fois de comparer et de calculer la quantité de vidéo de longueur arbitraire qui sera compressée. Parce que aujourd'hui nous envisageons la technologie du "vert", alors les graphiques seront appropriés -)

La vitesse d'encodage dépend directement de la qualité, c'est évident. Il convient de noter que la résolution de la lumière (appelons-la traditionnellement - SD) n'est pas un problème pour Badaboom - la vitesse d'encodage est 5,5 fois supérieure à la fréquence d'images vidéo d'origine (24 ips). Et même une vidéo 1080p lourde est convertie par le programme en temps réel. Il est à noter que la qualité de la vidéo finale est très proche de la vidéo originale, c'est-à-dire encode Badaboom de très, très haute qualité.

Mais généralement, la vidéo est dépassée à une résolution inférieure, voyons comment les choses se passent dans ce mode. Comme la résolution a été abaissée, le débit vidéo a également chuté. C'était 9500 kbps pour le fichier de sortie 1080p, 4100 kbps pour 720p et 2400 kbps pour 720x404. Le choix a été fait sur la base d'un rapport taille/qualité raisonnable.

Les commentaires sont superflus. Si vous passez de 720p à une qualité SD normale, il faudra environ 30 minutes pour transcoder un film de 2 heures. Et en même temps, la charge du processeur sera insignifiante, vous pourrez vaquer à vos occupations sans ressentir de gêne.

Mais que se passe-t-il si vous convertissez la vidéo dans un format pour un appareil mobile ? Pour cela, sélectionnez le profil iPhone (bitrate 1 Mbps, 480x320) et regardez la vitesse d'encodage :

Dois-je dire quelque chose ? Un film de deux heures en qualité iPhone normale est transcodé en moins de 15 minutes. La qualité HD est plus difficile, mais reste assez rapide. L'essentiel est que la qualité du matériel vidéo de sortie reste à un niveau assez élevé lorsqu'il est visualisé sur l'écran du téléphone.

En général, les impressions de Badaboom sont positives, la rapidité de travail plaît, l'interface est simple et directe. Toutes sortes de bugs dans les versions précédentes (j'utilisais encore la version bêta en 2008) ont été corrigés. À l'exception d'une chose - le chemin d'accès au fichier source, ainsi que le dossier dans lequel la vidéo terminée est enregistrée, ne doit pas contenir de lettres russes. Mais dans le contexte des mérites du programme, cet inconvénient est insignifiant.

Ensuite, nous aurons Super LoiLoScope... Pour la version habituelle, ils demandent 3 280 roubles, et pour la version tactile qui prend en charge le contrôle tactile dans Windows 7, ils demandent autant 4 440 roubles... Essayons de comprendre pourquoi le développeur veut ce genre d'argent et pourquoi l'éditeur vidéo a besoin d'un support multitouch. La dernière version a été utilisée - 1.8.3.3 .

Il est assez difficile de décrire l'interface du programme avec des mots, j'ai donc décidé de tourner une courte vidéo. Je dois dire tout de suite que, comme tous les convertisseurs vidéo pour CUDA, l'accélération GPU n'est prise en charge que pour la sortie vidéo en MPEG4 avec le codec h.264.

Lors de l'encodage, la charge du processeur est de 100%, mais cela ne provoque pas d'inconfort. Le navigateur et les autres applications non lourdes ne ralentissent pas.

Passons maintenant aux performances. Pour commencer, tout est identique à Badaboom - transcodage vidéo dans la même qualité.

Les résultats sont bien meilleurs que Badaboom. La qualité est également au top, la différence avec l'original ne se voit qu'en comparant les cadres par paires sous une loupe.

Wow, ici LoiloScope contourne Badaboom 2,5 fois. Dans le même temps, vous pouvez facilement couper et encoder une autre vidéo en parallèle, lire les actualités et même regarder un film, et même le FullHD peut être lu sans problème, bien que la charge du processeur soit maximale.

Essayons maintenant de faire une vidéo pour un appareil mobile, nous nommerons le profil de la même manière qu'il s'appelait dans Badaboom - iPhone (480x320, 1 Mbps):

Il n'y a pas d'erreur. Tout a été vérifié plusieurs fois, à chaque fois le résultat était le même. Très probablement, cela se produit pour la simple raison que le fichier SD est écrit avec un codec différent et dans un conteneur différent. Lors du transcodage, la vidéo est d'abord décodée, divisée en matrices d'une certaine taille et compressée. Le décodeur ASP utilisé en cas de xvid est plus lent que AVC (pour h.264) en décodage parallèle. Cependant, 192 ips est 8 fois plus rapide que la vitesse vidéo d'origine, une rafale de 23 minutes est compressée en moins de 4 minutes. La situation s'est répétée avec d'autres fichiers compressés en xvid / DivX.

LoiloScope n'a laissé que des impressions agréables sur moi-même - l'interface, malgré son caractère inhabituel, est pratique et fonctionnelle, et la vitesse de travail est au-delà des louanges. La fonctionnalité relativement médiocre est quelque peu frustrante, mais souvent avec une édition simple, il vous suffit d'ajuster légèrement les couleurs, de faire des transitions fluides, de superposer du texte, et LoiloScope fait un excellent travail avec cela. Le prix est également quelque peu effrayant - plus de 100 $ pour la version régulière est normal pour les pays étrangers, mais de tels chiffres nous semblent toujours un peu fous. Cependant, j'avoue que si, par exemple, j'avais souvent filmé et monté des vidéos personnelles, j'aurais peut-être pensé à acheter. En parallèle, j'ai d'ailleurs vérifié la possibilité d'éditer du contenu HD (ou plutôt AVCHD) directement depuis le caméscope sans le convertir au préalable dans un autre format, LoiloScope n'a révélé aucun problème avec les fichiers mts.

La nouvelle technologie est comme une espèce évolutive nouvellement apparue. Une créature étrange, contrairement aux nombreux anciens. Gênant à certains endroits, drôle à certains endroits. Et au premier abord, ses nouvelles qualités ne semblent en aucun cas adaptées à ce monde habitable et stable.

Cependant, un peu de temps passe, et il s'avère que le débutant court plus vite, saute plus haut et généralement plus fort. Et il mange plus de mouches que ses voisins rétrogrades. Et puis ces mêmes voisins commencent à comprendre qu'il ne vaut pas la peine de se quereller avec cet ancien maladroit. Mieux vaut être ami avec lui, et encore mieux organiser une symbiose. Vous regardez, et plus de mouches tomberont.

La technologie GPGPU (General-Purpose Graphics Processing Units) n'a longtemps existé que dans les calculs théoriques des universitaires intelligents. Sinon comment? Seuls les théoriciens sont capables de proposer de changer radicalement le processus de calcul qui s'est développé au fil des décennies en confiant le calcul de ses branches parallèles à une carte vidéo.

Le logo de la technologie CUDA nous rappelle qu'elle a grandi dans les profondeurs
Graphiques 3D.

Mais la technologie GPGPU n'allait pas prendre la poussière sur les pages des revues universitaires pendant longtemps. Ayant gonflé des plumes de ses meilleures qualités, elle a attiré l'attention des fabricants. C'est ainsi qu'est né CUDA - une implémentation de GPGPU sur les GPU NVIDIA GeForce.

Grâce à CUDA, les technologies GPGPU sont devenues courantes. Et maintenant, seul le plus myope et recouvert d'une épaisse couche de paresse, le développeur de systèmes de programmation ne déclare pas le support de son produit CUDA. IT-publications considérait comme un honneur de présenter les détails de la technologie dans de nombreux articles de vulgarisation scientifique, et les concurrents se sont urgemment assis sur des modèles et des compilateurs croisés pour développer quelque chose de similaire.

La reconnaissance publique n'est pas seulement le rêve des starlettes en herbe, mais aussi des nouvelles technologies émergentes. Et CUDA a de la chance. Elle est entendue, les gens parlent et écrivent à son sujet.

Ils écrivent juste comme s'ils continuaient à discuter de la GPGPU dans des revues scientifiques épaisses. Le lecteur est bombardé d'un tas de termes tels que "grille", "SIMD", "chaîne", "hôte", "texture et mémoire constante". Ils l'immergent au plus haut dans les schémas d'organisation des GPU nVidia, mènent les chemins sinueux des algorithmes parallèles et (le mouvement le plus fort) montrent de longues listes de code C. En conséquence, il s'avère qu'à l'entrée de l'article, nous avons un lecteur CUDA frais et désireux de comprendre, et à la sortie - le même lecteur, mais avec une tête enflée remplie d'un fouillis de faits, de diagrammes, de code, algorithmes et termes.

Pendant ce temps, le but de toute technologie est de nous rendre la vie plus facile. Et CUDA fait un excellent travail. Les résultats de son travail - c'est ce qui convaincra tout sceptique mieux qu'une centaine de schémas et d'algorithmes.

Loin de partout

CUDA est pris en charge par des supercalculateurs hautes performances
nVidia Tesla.

Et pourtant, avant de regarder les résultats des travaux de CUDA dans le domaine de la simplification de la vie de l'utilisateur moyen, il convient d'en comprendre toutes les limites. Comme un génie : tout désir, mais un. CUDA a aussi ses talons d'Achille. L'un d'eux est la limitation des plates-formes sur lesquelles elle peut travailler.

La liste des cartes vidéo fabriquées par nVidia qui prennent en charge CUDA est présentée dans une liste spéciale appelée CUDA Enabled Products. La liste est assez impressionnante, mais facilement classable. La prise en charge de CUDA n'est pas refusée :

Modèles NVidia GeForce 8e, 9e, 100e, 200e et 400e séries avec un minimum de 256 Mo de mémoire vidéo à bord. La prise en charge s'étend à la fois aux cartes de bureau et aux solutions mobiles.

La grande majorité des cartes graphiques de bureau et mobiles sont nVidia Quadro.

Toutes les solutions de la gamme de netbooks nvidia ION.

Solutions de superordinateur HPC (High Performance Computing) et nVidia Tesla hautes performances utilisées pour l'informatique personnelle ainsi que pour l'organisation de systèmes de cluster évolutifs.

Par conséquent, il vaut la peine de vérifier cette liste de favoris avant d'utiliser des produits logiciels basés sur CUDA.

En plus de la carte vidéo elle-même, un pilote correspondant est requis pour prendre en charge CUDA. C'est lui qui fait le lien entre le processeur central et le processeur graphique, agissant comme une sorte d'interface logicielle pour accéder au code du programme et aux données de la trésorerie multicœur du GPU. Pour être sûr de ne pas se tromper, nVidia vous recommande de visiter la page des pilotes pour la version la plus récente.

... mais le processus lui-même

Comment fonctionne CUDA ? Comment expliquer le processus complexe du calcul parallèle sur une architecture matérielle particulière du GPU pour ne pas plonger le lecteur dans l'abîme de termes précis ?

Vous pouvez essayer de le faire en imaginant comment le processeur central exécute le programme en symbiose avec le processeur graphique.

L'unité centrale de traitement architecturale (CPU) et son frère graphique (GPU) sont structurés différemment. Si nous faisons une analogie avec le monde de l'industrie automobile, alors le CPU est un break, l'un de ceux que l'on appelle une "grange". Cela ressemble à une voiture de tourisme, mais en même temps (du point de vue des développeurs) "un Suisse, et un faucheur, et un joueur sur un tuyau". Joue à la fois le rôle d'un petit camion, d'un bus et d'un hayon hypertrophié. Chariot, en somme. Il a peu de cylindres de base, mais ils "tirent" presque toutes les tâches, et l'impressionnante mémoire cache est capable d'accueillir un tas de données.

Le GPU, quant à lui, est une voiture de sport. Une seule fonction : amener le pilote à la ligne d'arrivée le plus rapidement possible. Par conséquent, pas de grand coffre à mémoire, pas de sièges supplémentaires. Mais le nombre de cylindres de noyau est des centaines de fois supérieur à celui d'un processeur.

Grâce à CUDA, les développeurs de programmes GPGPU n'ont pas besoin de se plonger dans la complexité du programme
rationnement pour les moteurs graphiques tels que DirectX et OpenGL

Contrairement à un processeur central capable de résoudre n'importe quelle tâche, y compris graphique, mais avec des performances moyennes, le processeur graphique est adapté à une solution à grande vitesse d'un problème : convertir des tas de polygones en entrée en un tas de pixels en sortie. De plus, ce problème peut être résolu en parallèle sur des centaines de cœurs de calcul relativement simples dans le GPU.

Alors, que peut être un tandem d'un break et d'une voiture de sport ? Le travail de CUDA ressemble à ceci : un programme est exécuté sur le CPU jusqu'à ce qu'un morceau de code y apparaisse qui peut être exécuté en parallèle. Ensuite, au lieu de s'exécuter lentement sur deux (voire huit) cœurs du processeur le plus cool, il est transféré sur des centaines de cœurs GPU. Dans le même temps, le temps d'exécution de cette section est considérablement réduit, ce qui signifie que le temps d'exécution de l'ensemble du programme est également réduit.

Rien ne change technologiquement pour le programmeur. Le code CUDA est écrit en langage C. Plus précisément, dans son dialecte spécial "C avec ruisseaux" (C avec ruisseaux). Développée à Stanford, cette extension en C s'appelle Brook. L'interface de transmission du code Brook au GPU est le pilote d'une carte vidéo prenant en charge CUDA. Il organise l'ensemble du processus de traitement de cette section du programme de sorte que pour le programmeur GPU, il ressemble à un coprocesseur du CPU. Tout comme l'utilisation d'un coprocesseur mathématique au début des ordinateurs personnels. Avec l'avènement de Brook, des cartes vidéo et des pilotes compatibles CUDA, tout programmeur peut accéder au GPU dans ses programmes. Mais avant, ce chamanisme appartenait à un cercle restreint de l'élite, perfectionnant depuis des années les techniques de programmation pour les moteurs graphiques DirectX ou OpenGL.

Dans le tonneau de ce miel prétentieux - les louanges de CUDA - vous devriez mettre une mouche dans la pommade, c'est-à-dire des restrictions. Toutes les tâches qui doivent être programmées ne peuvent pas être résolues avec CUDA. Cela ne fonctionnera pas pour accélérer la résolution des tâches de bureau de routine, mais confiez à CUDA le soin de calculer le comportement de milliers de combattants du même type dans World of Warcraft - s'il vous plaît. Mais c'est une tâche incontrôlable. Regardons des exemples de ce que CUDA résout déjà très efficacement.

uvres justes

CUDA est une technologie très pragmatique. Avec son support de ses cartes graphiques, nVidia espérait à juste titre que la bannière CUDA serait reprise par une multitude de passionnés tant au niveau universitaire que professionnel. Et ainsi c'est arrivé. Les projets basés sur CUDA vivent et profitent.

NVIDIA PhysX

En annonçant le prochain chef-d'œuvre du jeu, les fabricants insistent souvent sur son réalisme 3D. Mais peu importe à quel point le monde du jeu en 3D est réel, si les lois élémentaires de la physique, telles que la gravité, le frottement, l'hydrodynamique, sont mises en œuvre de manière incorrecte, la fausseté sera ressentie instantanément.

L'une des caractéristiques du moteur physique NVIDIA PhysX est la manipulation réaliste des tissus.

La mise en œuvre d'algorithmes pour la simulation informatique des lois physiques de base est très laborieuse. Les entreprises les plus connues dans ce domaine sont la société irlandaise Havok avec sa physique multiplateforme Havok Physics et la californienne Ageia - l'ancêtre du premier processeur physique au monde (PPU - Physics Processing Unit) et le moteur physique correspondant PhysX. Le premier d'entre eux, bien qu'acquis par Intel, travaille désormais activement à l'optimisation du moteur Havok pour les cartes vidéo ATI et les processeurs AMD. Mais Ageia avec son moteur PhysX est devenu une partie de nVidia. Dans le même temps, nVidia a résolu la tâche assez difficile d'adapter PhysX à la technologie CUDA.

Cela est devenu possible grâce aux statistiques. Il a été statistiquement prouvé que peu importe la complexité du rendu effectué par le GPU, certains de ses cœurs sont toujours inactifs. C'est sur ces cores que fonctionne le moteur PhysX.

Grâce à CUDA, la part du lion des calculs liés à la physique du monde du jeu a été réalisée sur une carte vidéo. La puissance libérée du processeur central a été utilisée pour résoudre d'autres problèmes de gameplay. Le résultat ne s'est pas fait attendre. Selon les experts, le gain de performances du gameplay basé sur PhysX sur CUDA a augmenté d'au moins un ordre de grandeur. La probabilité de la mise en œuvre des lois physiques a également augmenté. CUDA s'occupe du calcul de routine de la mise en œuvre du frottement, de la gravitation et d'autres choses auxquelles nous sommes habitués pour les objets multidimensionnels. Désormais, non seulement les héros et leur technique s'intègrent parfaitement dans les lois du monde physique auquel nous sommes habitués, mais aussi la poussière, le brouillard, les ondes de choc, les flammes et l'eau.

Version CUDA de NVIDIA Texture Tools 2

Vous aimez les objets réalistes dans les jeux modernes ? Merci aux développeurs de textures. Mais plus il y a de réalité dans la texture, plus son volume est grand. Plus il occupe une mémoire précieuse. Pour éviter cela, les textures sont précompressées et décompressées dynamiquement selon les besoins. Et compresser et décompresser demande beaucoup de calculs. Pour travailler avec des textures, nVidia a publié le package NVIDIA Texture Tools. Il prend en charge une compression et une décompression efficaces des textures DirectX (le soi-disant format RFC). La deuxième version de ce package prend en charge les algorithmes de compression BC4 et BC5, implémentés dans la technologie DirectX 11. Mais l'essentiel est que NVIDIA Texture Tools 2 prenne en charge CUDA. Selon nVidia, cela donne une amélioration de 12 fois des performances pour la compression et la décompression des textures. Cela signifie que les images du gameplay se chargeront plus rapidement et raviront le joueur par leur réalisme.

Le package NVIDIA Texture Tools 2 est conçu pour fonctionner avec CUDA. Les gains de performances pour la compression et la décompression des textures sont évidents.

L'utilisation de CUDA peut améliorer considérablement l'efficacité du suivi vidéo.

Traitement de flux vidéo en temps réel

Quoi qu'on en dise, mais le monde actuel, du point de vue de l'écoute, est beaucoup plus proche du monde du Big Brother d'Orwell qu'il n'y paraît. Les regards des caméras vidéo sont ressentis à la fois par les automobilistes et les visiteurs des lieux publics.

Des fleuves pleins d'informations vidéo se jettent dans les centres de son traitement et ... se heurtent à un lien étroit - une personne. C'est lui qui, dans la plupart des cas, est la dernière instance à surveiller le monde de la vidéo. De plus, l'autorité n'est pas la plus efficace. Clignote, est distrait et essaie de s'endormir.

Grâce à CUDA, il est devenu possible de mettre en œuvre des algorithmes de suivi simultané de plusieurs objets dans un flux vidéo. Dans ce cas, le processus se déroule en temps réel et la vidéo est à part entière à 30 ips. Par rapport à la mise en œuvre d'un tel algorithme sur les processeurs multicœurs modernes, CUDA donne une augmentation de deux ou trois fois des performances, et c'est beaucoup, voyez-vous.

Conversion vidéo, filtrage audio

Badaboom Video Converter est la première hirondelle à utiliser CUDA pour accélérer les conversions.

C'est agréable de regarder la nouvelle vidéo de location en qualité FullHD et sur grand écran. Mais vous ne pouvez pas emporter un grand écran avec vous sur la route, et le codec vidéo FullHD aura le hoquet sur le processeur basse consommation d'un gadget mobile. La conversion vient à la rescousse. Mais la plupart de ceux qui l'ont rencontré dans la pratique se plaignent du long temps de conversion. Cela se comprend, le processus est routinier, adapté à la parallélisation, et son exécution sur le CPU n'est pas très optimale.

Mais CUDA y fait face avec brio. La première hirondelle est le convertisseur Badaboom d'Elevental. Les développeurs de Badaboom ne se sont pas trompés en choisissant CUDA. Les tests montrent qu'un film standard d'une heure et demie est converti au format iPhone/iPod Touch en moins de vingt minutes. Et ce malgré le fait qu'en utilisant uniquement le CPU, ce processus prend plus d'une heure.

Aide CUDA et les mélomanes professionnels. N'importe lequel d'entre eux abandonnera mon royaume pour un crossover FIR efficace - un ensemble de filtres divisant le spectre audio en plusieurs bandes. Ce processus est très laborieux et avec un grand volume de matériel audio, l'ingénieur du son va "fumer" pendant plusieurs heures. La mise en œuvre du croisement FIR basé sur CUDA accélère son fonctionnement des centaines de fois.

Avenir CUDA

En faisant de la technologie GPGPU une réalité, CUDA ne va pas se reposer sur ses lauriers. Comme partout, le principe de réflexion fonctionne dans CUDA : désormais, non seulement l'architecture des processeurs vidéo nVidia affecte le développement des versions du SDK CUDA, mais la technologie CUDA elle-même oblige nVidia à revoir l'architecture de ses puces. Un exemple d'une telle réflexion est la plate-forme nVidia ION. Sa deuxième version est spécialement optimisée pour résoudre les tâches CUDA. Cela signifie que même dans des solutions matérielles relativement peu coûteuses, les consommateurs bénéficieront de toute la puissance et l'éclat de CUDA.

Revenons à l'histoire - en 2003, lorsque Intel et AMD ont participé à la course commune au processeur le plus puissant. En quelques années seulement, cette course s'est traduite par une augmentation significative des vitesses d'horloge, notamment depuis la sortie de l'Intel Pentium 4.

Mais la course approchait rapidement de sa limite. Après une vague d'augmentations énormes des vitesses d'horloge (entre 2001 et 2003, la vitesse d'horloge du Pentium 4 a doublé de 1,5 à 3 GHz), les utilisateurs ont dû se contenter de dixièmes de gigahertz, que les fabricants ont pu évincer (de 2003 à 2005, les vitesses d'horloge sont passées de seulement 3 à 3, 8 GHz).

Même les architectures optimisées pour des vitesses d'horloge élevées, comme Prescott, ont commencé à rencontrer des difficultés, et cette fois pas seulement celles de production. Les fabricants de puces se sont simplement heurtés aux lois de la physique. Certains analystes ont même prédit que la loi de Moore cesserait de fonctionner. Mais cela ne s'est pas produit. Le sens originel de la loi est souvent déformé, mais il concerne le nombre de transistors à la surface d'un noyau de silicium. Pendant longtemps, l'augmentation du nombre de transistors dans le CPU s'est accompagnée d'une augmentation correspondante des performances - ce qui a conduit à une distorsion du sens. Mais ensuite, la situation s'est compliquée. Les concepteurs de l'architecture CPU se sont approchés de la loi de réduction de gain : le nombre de transistors à ajouter pour l'augmentation souhaitée des performances est devenu de plus en plus important, conduisant à une impasse.

Alors que les fabricants de processeurs s'arrachaient les derniers cheveux en essayant de trouver des solutions à leurs problèmes, les fabricants de GPU continuaient de bénéficier remarquablement des avantages de la loi de Moore.

Pourquoi n'ont-ils pas atteint la même impasse que les architectes CPU ? La raison est très simple : les processeurs sont conçus pour des performances maximales sur un flux d'instructions qui traitent différentes données (entiers et nombres à virgule flottante), effectuent un accès aléatoire à la mémoire, etc. Jusqu'à présent, les développeurs essayaient de fournir plus de parallélisme d'instructions, c'est-à-dire d'exécuter autant d'instructions en parallèle que possible. Ainsi, par exemple, avec le Pentium, l'exécution superscalaire est apparue, alors que dans certaines conditions, il était possible d'exécuter deux instructions par cycle. Le Pentium Pro a reçu une exécution désordonnée d'instructions, ce qui a permis d'optimiser le travail des unités de calcul. Le problème est que l'exécution parallèle d'un flux séquentiel d'instructions a des limites évidentes, donc une augmentation aveugle du nombre d'unités de calcul ne donne aucun avantage, car elles resteront inactives la plupart du temps.

En revanche, le travail du GPU est relativement simple. Elle consiste à accepter un groupe de polygones d'un côté et à générer un groupe de pixels de l'autre. Les polygones et les pixels sont indépendants les uns des autres, ils peuvent donc être traités en parallèle. Ainsi, dans le GPU, une grande partie du cristal peut être séparée en unités de calcul, qui, contrairement au CPU, seront effectivement utilisées.

Cliquez sur l'image pour agrandir.

Le GPU diffère du CPU non seulement en cela. L'accès à la mémoire dans le GPU est très lié - si un texel est lu, alors après quelques cycles d'horloge, le texel adjacent sera lu ; lorsqu'un pixel est enregistré, après quelques cycles d'horloge, le pixel adjacent sera enregistré. En organisant judicieusement la mémoire, vous pouvez obtenir des performances proches de la bande passante théorique. Cela signifie que le GPU, contrairement au CPU, ne nécessite pas un énorme cache, car son rôle est d'accélérer les opérations de texturation. Il suffit de quelques kilo-octets contenant plusieurs texels utilisés dans les filtres bilinéaires et trilinéaires.

Cliquez sur l'image pour agrandir.

Vive GeForce FX !

Les deux mondes sont restés longtemps séparés. Nous avons utilisé un processeur (ou même plusieurs processeurs) pour les tâches bureautiques et les applications Internet, et les GPU n'étaient bons que pour accélérer le rendu. Mais une chose a tout changé : à savoir, l'avènement des GPU programmables. Au début, les unités centrales n'avaient rien à craindre. Les premiers GPU dits programmables (NV20 et R200) ​​n'étaient guère une menace. Le nombre d'instructions dans le programme est resté limité à environ 10, elles fonctionnaient sur des types de données très exotiques tels que des nombres à virgule fixe de 9 ou 12 bits.

Cliquez sur l'image pour agrandir.

Mais la loi de Moore a de nouveau montré son meilleur côté. L'augmentation du nombre de transistors a non seulement augmenté le nombre d'unités de calcul, mais a également amélioré leur flexibilité. L'apparition du NV30 peut être considérée comme une avancée significative pour plusieurs raisons. Bien sûr, les joueurs n'aimaient pas vraiment les cartes NV30, mais les nouveaux GPU ont commencé à s'appuyer sur deux fonctionnalités conçues pour changer la perception des GPU non seulement en tant qu'accélérateurs graphiques.

• Prise en charge des calculs à virgule flottante simple précision (même s'il n'est pas conforme à la norme IEEE754) ;
• support pour plus d'un millier d'instructions.

On a donc toutes les conditions pour attirer des chercheurs pionniers qui sont toujours à la recherche de puissance de calcul supplémentaire.

L'idée d'utiliser des accélérateurs graphiques pour des calculs mathématiques n'est pas nouvelle. Les premières tentatives ont été faites dans les années 90 du siècle dernier. Bien sûr, ils étaient très primitifs - limités, pour la plupart, à l'utilisation de certaines fonctionnalités matérielles, telles que la rastérisation et les tampons Z pour accélérer des tâches telles que la recherche d'un itinéraire ou le rendu. Diagrammes de Voronoï .

Cliquez sur l'image pour agrandir.

En 2003, avec l'avènement des shaders évolués, une nouvelle référence a été atteinte - cette fois en effectuant des calculs matriciels. C'est l'année qu'une section entière de SIGGRAPH ("Calculs sur GPU") a été consacrée au nouveau domaine de l'informatique. Cette première initiative s'appelait GPGPU (General-Purpose calcul on GPU). Et le premier tournant a été l'apparence.

Pour comprendre le rôle de BrookGPU, vous devez comprendre comment tout s'est passé avant qu'il n'apparaisse. Le seul moyen d'obtenir des ressources GPU en 2003 était d'utiliser l'une des deux API graphiques, Direct3D ou OpenGL. Par conséquent, les développeurs qui voulaient obtenir la puissance du GPU pour leur calcul devaient s'appuyer sur les deux API mentionnées. Le problème est qu'ils n'étaient pas toujours experts en programmation de cartes vidéo, ce qui rendait l'accès à la technologie très difficile. Alors que les programmeurs 3D opèrent avec des shaders, des textures et des fragments, les spécialistes du domaine de la programmation parallèle s'appuient sur des threads, des cœurs, des scatter, etc. Il fallait donc d'abord établir des analogies entre les deux mondes.

• Flux est un flux d'éléments du même type, dans le GPU il peut être représenté par une texture. Fondamentalement, dans la programmation classique, il existe un analogue tel qu'un tableau.
• Noyau- une fonction qui sera appliquée indépendamment à chaque élément du flux ; est l'équivalent d'un pixel shader. En programmation classique, l'analogie d'une boucle peut être faite - elle s'applique à un grand nombre d'éléments.
• Pour lire les résultats de l'application d'un noyau à un flux, une texture doit être créée. Il n'y a pas d'équivalent sur le CPU, car il y a un accès complet à la mémoire.
• L'emplacement en mémoire où l'enregistrement sera effectué (dans les opérations scatter / scatter) est contrôlé via le vertex shader, car le pixel shader ne peut pas changer les coordonnées du pixel traité.

Comme vous pouvez le voir, même en tenant compte des analogies données, la tâche ne semble pas simple. Et Brook est venu à la rescousse. Ce nom fait référence aux extensions du langage C ("C with streams", "C with streams"), comme les appelaient les développeurs de Stanford. À la base, la tâche de Brook était de cacher au programmeur tous les composants de l'API 3D, ce qui permettait de présenter le GPU comme un coprocesseur pour des calculs parallèles. Pour ce faire, le compilateur Brook a traité un fichier .br avec du code et des extensions C ++, puis a généré du code C ++ lié à une bibliothèque prenant en charge différentes sorties (DirectX, OpenGL ARB, OpenGL NV3x, x86).

Cliquez sur l'image pour agrandir.

Brook a plusieurs mérites, dont le premier est de faire sortir le GPGPU de l'ombre afin que le grand public puisse se familiariser avec la technologie. Cependant, après l'annonce du projet, un certain nombre de sites informatiques étaient trop optimistes que la sortie de Brook jette le doute sur l'existence des CPU, qui seront bientôt remplacés par des GPU plus puissants. Mais, comme on peut le voir, cela ne s'est pas produit après cinq ans. Pour être honnête, nous pensons que cela n'arrivera jamais. D'autre part, en regardant l'évolution réussie des CPU, de plus en plus orientés vers le parallélisme (plus de cœurs, technologie multithreading SMT, expansion des blocs SIMD), ainsi que des GPU, qui, au contraire, deviennent de plus en plus universels (prise en charge des calculs en virgule flottante simple précision, calcul d'entiers, prise en charge de la double précision), il semble que GPU et CPU vont bientôt fusionner. Que se passera-t-il alors ? Les GPU seront-ils engloutis par le CPU, comme cela s'est produit avec les coprocesseurs mathématiques ? Tout à fait possible. Intel et AMD travaillent actuellement sur des projets similaires. Mais il y a encore beaucoup de choses qui peuvent changer.

Mais revenons à notre sujet. L'avantage de Brook était de vulgariser le concept GPGPU, il simplifiait considérablement l'accès aux ressources GPU, ce qui permettait à de plus en plus d'utilisateurs de maîtriser le nouveau modèle de programmation. D'un autre côté, malgré toutes les qualités de Brook, il y avait encore un long chemin à parcourir avant que les ressources GPU puissent être utilisées pour le calcul.

L'un des problèmes est lié aux différents niveaux d'abstraction, et aussi, en particulier, à la charge supplémentaire excessive créée par l'API 3D, qui peut être assez sensible. Mais plus grave est le problème de compatibilité sur lequel les développeurs de Brook n'ont rien pu faire. La concurrence est féroce entre les fabricants de GPU, ils optimisent donc souvent leurs pilotes. Si ces optimisations sont bonnes pour la plupart des joueurs, elles pourraient mettre fin à la compatibilité Brook dans un instant. Par conséquent, il est difficile d'imaginer utiliser cette API dans du code de production qui fonctionnera quelque part. Et pendant longtemps, Brook est resté le lot des chercheurs et programmeurs amateurs.

Cependant, le succès de Brook a suffi à attirer l'attention d'ATI et de nVidia, et ils ont développé un intérêt pour une telle initiative, car elle pourrait élargir le marché, ouvrant un nouveau secteur important pour les entreprises.

Les chercheurs initialement impliqués dans le projet Brook ont ​​rapidement rejoint les équipes de développement à Santa Clara pour présenter une stratégie globale de développement d'un nouveau marché. L'idée était de créer une combinaison de matériel et de logiciels adaptés aux tâches GPGPU. Étant donné que les développeurs nVidia connaissent tous les secrets de leurs GPU, on ne pouvait pas se fier à l'API graphique, mais plutôt à communiquer avec le GPU via un pilote. Bien que, bien sûr, cela a ses propres problèmes. Ainsi, l'équipe de développement CUDA (Compute Unified Device Architecture) a créé un ensemble de couches logicielles pour travailler avec le GPU.

Cliquez sur l'image pour agrandir.

Comme vous pouvez le voir dans le diagramme, CUDA fournit deux API.

• API de haut niveau : API d'exécution CUDA ;
• API de bas niveau : API de pilote CUDA.

Étant donné que l'API de haut niveau est implémentée au-dessus de l'API de bas niveau, chaque appel de fonction d'exécution est décomposé en instructions plus simples que l'API du pilote traite. Notez que les deux API s'excluent mutuellement : un programmeur peut utiliser l'une ou l'autre API, mais mélanger les appels aux fonctions des deux API ne fonctionnera pas. En général, le terme « API de haut niveau » est relatif. Même l'API Runtime est telle que beaucoup la considéreraient comme de bas niveau ; cependant, il fournit des fonctions assez pratiques pour l'initialisation ou la gestion du contexte. Mais ne vous attendez pas à un niveau d'abstraction particulièrement élevé - vous devez tout de même avoir une bonne connaissance des GPU nVidia et de leur fonctionnement.

L'API du pilote est encore plus difficile à utiliser ; il faut plus d'efforts pour exécuter le traitement GPU. D'un autre côté, l'API de bas niveau est plus flexible, donnant au programmeur un contrôle supplémentaire si nécessaire. Deux API sont capables de fonctionner avec des ressources OpenGL ou Direct3D (seulement la neuvième version pour aujourd'hui). Les avantages de cette fonctionnalité sont évidents - CUDA peut être utilisé pour créer des ressources (géométrie, textures procédurales, etc.) qui peuvent être transmises à l'API graphique, ou, à l'inverse, vous pouvez faire en sorte que l'API 3D envoie les résultats du rendu à la CUDA programme, qui, à son tour, effectuera le post-traitement. Les exemples de ces interactions sont nombreux, et l'avantage est que les ressources continuent d'être stockées dans la mémoire du GPU sans avoir à passer sur le bus PCI Express, ce qui reste un goulot d'étranglement.

Cependant, il faut noter que le partage des ressources en mémoire vidéo n'est pas toujours parfait et peut conduire à quelques « casse-tête ». Par exemple, lors du changement de résolution ou de profondeur de couleur, les données graphiques sont prioritaires. Par conséquent, si vous avez besoin d'augmenter les ressources dans le framebuffer, le pilote le fera sans problème au détriment des ressources des applications CUDA, qui "décolleront" simplement avec une erreur. Pas très élégant, bien sûr, mais cela ne devrait pas arriver très souvent. Et pendant que nous commencions à parler des inconvénients : si vous souhaitez utiliser plusieurs GPU pour les applications CUDA, vous devez d'abord désactiver le mode SLI, sinon les applications CUDA ne pourront "voir" qu'un seul GPU.

Enfin, la troisième couche de logiciels est donnée aux bibliothèques - deux, pour être précis.

• CUBLAS, où se trouvent les blocs nécessaires au calcul de l'algèbre linéaire sur le GPU ;
• CUFFT, qui prend en charge le calcul des transformées de Fourier, un algorithme largement utilisé dans le traitement du signal.

Avant de plonger dans CUDA, permettez-moi de définir un certain nombre de termes éparpillés dans la documentation nVidia. L'entreprise a choisi une terminologie très spécifique à laquelle il est difficile de s'habituer. Tout d'abord, notez que fil dans CUDA est loin d'avoir la même signification que thread CPU, ni l'équivalent de thread dans nos articles sur les GPU. Le thread GPU dans ce cas est l'ensemble de données sous-jacent qui doit être traité. Contrairement aux threads CPU, les threads CUDA sont très "légers", c'est-à-dire que la commutation de contexte entre deux threads n'est en aucun cas une opération gourmande en ressources.

Le deuxième terme souvent trouvé dans la documentation CUDA est chaîne... Il n'y a pas de confusion ici, puisqu'il n'y a pas d'analogue en russe (sauf si vous êtes un fan de Start Trek ou du jeu Warhammer). En fait, le terme est tiré de l'industrie textile, où le fil de trame est tiré à travers le fil de chaîne, qui est étiré sur la machine. Warp dans CUDA est un groupe de 32 threads et représente la quantité minimale de données traitées par la méthode SIMD dans les multiprocesseurs CUDA.

Mais ce "grain" n'est pas toujours pratique pour le programmeur. Par conséquent, dans CUDA, au lieu de travailler directement avec les chaînes, vous pouvez travailler avec blocs / bloc contenant de 64 à 512 threads.

Enfin, ces blocs se rejoignent en grilles / grille... L'avantage de ce regroupement est que le nombre de blocs traités simultanément par le GPU est étroitement lié aux ressources matérielles, comme nous le verrons plus loin. Le regroupement de blocs dans des grilles vous permet de vous affranchir complètement de cette limitation et d'appliquer le noyau/noyau à plusieurs threads en un seul appel, sans penser aux ressources fixes. Les bibliothèques CUDA sont responsables de tout cela. De plus, un tel modèle évolue bien. Si le GPU a peu de ressources, il exécutera les blocs de manière séquentielle. Si le nombre de processeurs de calcul est grand, alors les blocs peuvent être exécutés en parallèle. C'est-à-dire que le même code peut s'exécuter à la fois sur les GPU d'entrée de gamme et sur les modèles haut de gamme et même futurs.

Il y a quelques autres termes dans l'API CUDA qui font référence au processeur ( hôte / hôte) et GPU ( appareil / appareil). Si cette petite introduction ne vous a pas fait peur, alors il est temps de regarder de plus près CUDA.

Si vous lisez régulièrement le Guide du matériel de Tom, alors l'architecture des derniers GPU de nVidia vous est familière. Sinon, nous vous recommandons de consulter l'article " nVidia GeForce GTX 260 et 280 : la prochaine génération de cartes graphiques"En termes de CUDA, nVidia présente l'architecture d'une manière légèrement différente, montrant certains détails qui étaient auparavant cachés.

Comme vous pouvez le voir sur l'illustration ci-dessus, le noyau de shader nVidia est composé de plusieurs clusters de processeurs de texture. (Groupe de processeurs de texture, TPC)... La 8800 GTX, par exemple, utilisait huit clusters, la 8800 GTS six, et ainsi de suite. Chaque cluster, en fait, se compose d'une unité de texture et de deux multiprocesseur de diffusion en continu... Ces derniers comprennent le début du pipeline (front end), qui lit et décode les instructions, ainsi que leur envoi pour exécution, et la fin du pipeline (back end), qui se compose de huit appareils informatiques et de deux appareils super fonctionnels. . SFU (Unité Super Fonctionnelle) où les instructions sont exécutées selon le principe SIMD, c'est-à-dire qu'une instruction s'applique à tous les threads de la chaîne. nVidia appelle cette façon de faire SIMT(une seule instruction plusieurs threads, une instruction, plusieurs threads). Il est important de noter que la fin du pipeline fonctionne à deux fois la fréquence de son début. En pratique, cela signifie que cette partie semble deux fois plus "large" qu'elle ne l'est réellement (c'est-à-dire comme un bloc SIMD à 16 canaux au lieu d'un bloc à huit canaux). Les multiprocesseurs de streaming fonctionnent comme suit : à chaque cycle d'horloge, le début du pipeline sélectionne une chaîne prête à être exécutée et commence à exécuter l'instruction. Il faudrait quatre cycles d'horloge pour que l'instruction s'applique aux 32 threads du warp, mais comme elle s'exécute à deux fois la fréquence du démarrage, cela ne prend que deux cycles d'horloge (en termes de démarrage du pipeline). Par conséquent, pour que le début du pipeline ne ralentisse pas un cycle et que le matériel soit chargé autant que possible, dans le cas idéal, vous pouvez alterner les instructions à chaque cycle - une instruction classique en un cycle et une instruction pour SFU - en un autre.

Chaque multiprocesseur possède un ensemble spécifique de ressources qui valent la peine d'être compris. Il y a une petite zone de mémoire appelée "La memoire partagée", 16 Ko par multiprocesseur. Il ne s'agit en aucun cas d'une mémoire cache : le programmeur peut l'utiliser à sa discrétion. C'est-à-dire que nous avons devant nous quelque chose de proche du magasin local des processeurs SPU on Cell. Ce détail est curieux car il souligne que CUDA est une combinaison de technologies logicielles et matérielles. Cette zone de mémoire n'est pas utilisée pour les pixel shaders, ce que nVidia souligne avec humour "nous n'aimons pas que les pixels se parlent".

Cette zone mémoire ouvre la possibilité d'échanger des informations entre les threads. en un seul bloc... Il est important de souligner cette limitation : tous les threads d'un bloc sont assurés d'être exécutés par un seul multiprocesseur. Au contraire, la liaison de blocs à différents multiprocesseurs n'est pas du tout stipulée et deux threads de blocs différents ne peuvent pas échanger d'informations entre eux lors de l'exécution. Autrement dit, l'utilisation de la mémoire partagée n'est pas facile. Cependant, la mémoire partagée reste justifiée, sauf dans les cas où plusieurs threads tentent d'accéder à la même banque mémoire, provoquant un conflit. Dans d'autres situations, l'accès à la mémoire partagée est aussi rapide qu'aux registres.

La mémoire partagée n'est pas la seule mémoire à laquelle les multiprocesseurs peuvent accéder. Ils peuvent utiliser de la mémoire vidéo, mais avec une bande passante plus faible et une latence plus élevée. Ainsi, afin de réduire la fréquence d'accès à cette mémoire, nVidia a équipé les multiprocesseurs d'un cache (environ 8 Ko par multiprocesseur) stockant les constantes et les textures.

Le multiprocesseur a 8192 registres, qui sont communs à tous les flux de tous les blocs actifs sur le multiprocesseur. Le nombre de blocs actifs par multiprocesseur ne peut pas dépasser huit et le nombre de warps actifs est limité à 24 (768 threads). Par conséquent, la 8800 GTX peut gérer jusqu'à 12 288 threads à la fois. Toutes ces limitations méritent d'être mentionnées car elles permettent d'optimiser l'algorithme en fonction des ressources disponibles.

Optimiser un programme CUDA consiste donc à obtenir un équilibre optimal entre le nombre de blocs et leur taille. Plus de threads par bloc sera utile pour réduire la latence de la mémoire, mais le nombre de registres disponibles par thread est également réduit. De plus, un bloc de 512 threads sera inefficace, puisqu'un seul bloc peut être actif sur un multiprocesseur, ce qui entraînera la perte de 256 threads. Par conséquent, nVidia recommande d'utiliser des blocs de 128 ou 256 threads, ce qui offre le meilleur compromis entre une latence plus faible et le nombre de registres pour la plupart des noyaux/noyau.

D'un point de vue programmatique, CUDA se compose d'un ensemble d'extensions C, rappelant BrookGPU, ainsi que de plusieurs appels d'API spécifiques. Les extensions incluent des spécificateurs de type liés aux fonctions et aux variables. Il est important de se souvenir du mot clé __global__, qui, étant donné avant la fonction, montre que cette dernière fait référence au noyau/noyau - cette fonction sera appelée par le CPU, et elle sera exécutée sur le GPU. Préfixe __dispositif__ indique que la fonction sera exécutée sur le GPU (que d'ailleurs CUDA appelle "device / device") mais elle ne peut être appelée que depuis le GPU (c'est-à-dire depuis une autre fonction __device__ ou depuis la fonction __global__). Enfin, le préfixe __hôte__ facultatif, il désigne une fonction appelée par le CPU et exécutée par le CPU - en d'autres termes, une fonction normale.

Il existe un certain nombre de limitations avec les fonctions __device__ et __global__ : elles ne peuvent pas être récursives (c'est-à-dire s'appeler elles-mêmes) et elles ne peuvent pas avoir un nombre variable d'arguments. Enfin, étant donné que les fonctions __device__ sont situées dans l'espace mémoire du GPU, il est logique que vous ne puissiez pas obtenir leur adresse. Les variables ont également un certain nombre de qualificatifs qui indiquent la zone de mémoire où elles seront stockées. Variable préfixée __partagé__ signifie qu'il sera stocké dans la mémoire partagée du multiprocesseur de streaming. L'appel de fonction __global__ est légèrement différent. Le fait est que, lors de l'appel, vous devez définir la configuration d'exécution - plus précisément, la taille de la grille / grille à laquelle le noyau / noyau sera appliqué, ainsi que la taille de chaque bloc. Prenez, par exemple, le noyau avec la signature suivante.

__global__ void Func (paramètre float *);

Il s'appellera comme

Fonction<<< Dg, Db >>> (paramètre);

où Dg est la taille de la grille et Db est la taille du bloc. Ces deux variables sont d'un nouveau type de vecteur introduit avec CUDA.

L'API CUDA contient des fonctions pour travailler avec la mémoire dans la VRAM : cudaMalloc pour allouer de la mémoire, cudaFree pour libérer et cudaMemcpy pour copier la mémoire entre la RAM et la VRAM et vice versa.

Nous terminerons ce survol d'une manière très intéressante de la compilation d'un programme CUDA : la compilation se fait en plusieurs étapes. Tout d'abord, le code CPU est récupéré et transmis au compilateur standard. Le code spécifique au GPU est d'abord converti dans le langage intermédiaire PTX. Il est similaire au langage assembleur et vous permet d'étudier votre code à la recherche d'inefficacités potentielles. Enfin, la dernière phase consiste à traduire le langage intermédiaire en instructions GPU spécifiques et à créer un fichier binaire.

Après avoir parcouru les documents nVidia, je veux juste essayer CUDA cette semaine. En effet, quoi de mieux que d'évaluer une API en créant son propre programme ? C'est à ce moment-là que la plupart des problèmes devraient faire surface, même si tout semble parfait sur le papier. De plus, la pratique montrera mieux à quel point vous avez bien compris tous les principes décrits dans la documentation CUDA.

Il est assez facile de se plonger dans un projet comme celui-ci. Aujourd'hui, un grand nombre d'outils gratuits mais de haute qualité sont disponibles en téléchargement. Pour notre test, nous avons utilisé Visual C++ Express 2005, qui a tout ce dont vous avez besoin. La partie la plus difficile a été de trouver un programme qui n'a pas pris des semaines à être porté sur le GPU, mais suffisamment amusant pour maintenir nos efforts en bonne forme. Enfin, nous avons choisi un morceau de code qui prend une heightmap et calcule la normalmap correspondante. Nous n'entrerons pas dans le détail de cette fonction, car elle n'est guère intéressante dans cet article. Bref, le programme s'occupe de la courbure des zones : pour chaque pixel de l'image initiale, on impose une matrice qui détermine la couleur du pixel résultant dans l'image générée par les pixels adjacents, selon une formule plus ou moins complexe. L'avantage de cette fonction est qu'elle est très facile à paralléliser, ce test démontre donc parfaitement les capacités de CUDA.

Un autre avantage est que nous avons déjà une implémentation sur le CPU, nous pouvons donc comparer son résultat avec la version CUDA - et ne pas réinventer la roue.

On répète encore une fois que le but du test était de se familiariser avec les utilitaires du SDK CUDA, et non de faire des tests comparatifs de versions pour CPU et GPU. Comme il s'agissait de notre première tentative de création d'un programme CUDA, nous avions peu d'espoir d'obtenir des performances élevées. Comme cette partie du code n'est pas critique, la version pour le CPU n'a pas été optimisée, donc une comparaison directe des résultats n'est guère intéressante.

Performance

Cependant, nous avons mesuré le temps d'exécution pour voir s'il y a un avantage à utiliser CUDA même avec l'implémentation la plus grossière, ou s'il faudra une pratique longue et fastidieuse pour obtenir une sorte de gain GPU. La machine de test a été prise dans notre laboratoire de développement - un ordinateur portable avec un processeur Core 2 Duo T5450 et une carte graphique GeForce 8600M GT sous Vista. C'est loin d'être un supercalculateur, mais les résultats sont assez intéressants, puisque le test n'est pas « affûté » pour le GPU. C'est toujours agréable de voir nVidia faire d'énormes gains sur des systèmes dotés de GPU monstrueux et d'une bande passante importante, mais en pratique, bon nombre des 70 millions de GPU compatibles CUDA sur le marché des PC d'aujourd'hui sont loin d'être aussi puissants, c'est pourquoi notre test a le droit de la vie.

Pour une image de 2 048 x 2 048 pixels, nous avons obtenu les résultats suivants.

• Processeur 1 thread : 1 419 ms ;
• CPU 2 threads : 749 ms ;
• CPU 4 threads : 593 ms
• Blocs GPU (8600M GT) de 256 threads : 109 ms ;
• Blocs GPU (8600M GT) de 128 threads : 94 ms ;
• Blocs GPU (8800 GTX) de 128 threads / 256 threads : 31 ms.

Plusieurs conclusions peuvent être tirées des résultats. Commençons par le fait que, malgré le discours sur la paresse évidente des programmeurs, nous avons modifié la version initiale du CPU pour plusieurs threads. Comme nous l'avons mentionné, le code est idéal pour cette situation - il suffit de diviser l'image initiale en autant de zones qu'il y a de flux. Veuillez noter que l'accélération de la transition d'un thread à deux sur notre processeur dual-core s'est avérée presque linéaire, ce qui indique également la nature parallèle du programme de test. De manière assez inattendue, la version à quatre threads s'est également avérée plus rapide, bien que cela soit très étrange sur notre processeur - au contraire, on pouvait s'attendre à une baisse d'efficacité en raison de la surcharge de gestion de threads supplémentaires. Comment expliquer ce résultat ? C'est difficile à dire, mais le planificateur de threads de Windows peut être à blâmer ; dans tous les cas, le résultat est reproductible. Avec des textures plus petites (512x512), le gain du découpage en threads n'était pas aussi prononcé (environ 35% contre 100%), et le comportement de la version à quatre threads était plus logique, sans augmentation par rapport à la version à deux threads. Le GPU était toujours plus rapide, mais pas aussi prononcé (8600M GT était trois fois plus rapide que la version dual-thread).

Cliquez sur l'image pour agrandir.

La deuxième observation importante est que même la mise en œuvre de GPU la plus lente s'est avérée presque six fois plus rapide que la version de CPU la plus performante. Pour le premier programme et la version non optimisée de l'algorithme, le résultat est très encourageant. Notez que nous avons obtenu un résultat nettement meilleur sur les petits blocs, bien que l'intuition puisse suggérer le contraire. L'explication est simple - notre programme utilise 14 registres par thread, et avec des blocs de 256 threads, 3 584 registres par bloc sont nécessaires et 768 threads sont nécessaires pour la charge complète du processeur, comme nous l'avons montré. Dans notre cas, il s'agit de trois blocs ou 10 572 registres. Mais le multiprocesseur n'a que 8192 registres, il ne peut donc garder que deux blocs actifs. En revanche, avec des blocs de 128 threads, nous avons besoin de 1792 registres par bloc ; si 8 192 est divisé par 1 792 et arrondi à l'entier le plus proche, alors nous obtenons quatre blocs. En pratique, le nombre de threads sera le même (512 par multiprocesseur, bien qu'en théorie 768 soient nécessaires pour la pleine charge), mais une augmentation du nombre de blocs donne au GPU l'avantage d'une flexibilité d'accès mémoire - lorsqu'une opération avec des latences élevées est en cours, vous pouvez lancer l'exécution des instructions d'un autre bloc, en attendant la réception des résultats. Quatre blocs réduisent nettement la latence, d'autant plus que notre programme utilise plusieurs accès mémoire.

Une analyse

Enfin, malgré ce que nous disions plus haut, nous n'avons pas pu résister à la tentation et avons fait tourner le programme sur la 8800 GTX, qui s'est avérée trois fois plus rapide que la 8600, quelle que soit la taille du bloc. On pourrait penser qu'en pratique, sur les architectures respectives, le résultat serait au moins quatre fois supérieur : 128 processeurs ALU / shader contre 32 et des vitesses d'horloge supérieures (1,35 GHz contre 950 MHz), mais cela n'a pas fonctionné. Le facteur limitant le plus probable était l'accès à la mémoire. Pour être plus précis, l'image initiale est accessible sous la forme d'un tableau CUDA multidimensionnel - un terme plutôt compliqué pour ce qui n'est rien de plus qu'une texture. Mais il y a plusieurs avantages.

• les accès bénéficient du cache de textures ;
• nous utilisons le mode wrapping, qui n'a pas besoin de gérer les bordures d'image, contrairement à la version CPU.

De plus, on peut profiter d'un filtrage "gratuit" avec un adressage normalisé entre et à la place de, mais dans notre cas cela n'est guère utile. Comme vous le savez, la 8600 possède 16 unités de texture contre 32 pour la 8800 GTX. Par conséquent, le rapport entre les deux architectures n'est que de deux pour un. Ajoutez à cela la différence de fréquences et nous obtenons le rapport (32 x 0,575) / (16 x 0,475) = 2,4 - proche du "trois pour un" que nous obtenons réellement. Cette théorie explique également pourquoi la taille des blocs ne change pas beaucoup sur le G80, puisque l'ALU repose toujours sur les blocs de texture.

Cliquez sur l'image pour agrandir.

Hormis des résultats prometteurs, notre première exposition à CUDA s'est très bien déroulée, compte tenu des conditions peu favorables choisies. Développer sur un ordinateur portable Vista, c'est utiliser le SDK CUDA 2.0, toujours en version bêta, avec le pilote 174.55, qui est également en version bêta. Malgré cela, nous ne pouvons signaler aucune mauvaise surprise - seulement des erreurs initiales lors du premier débogage, lorsque notre programme, encore assez "bogué", a essayé d'adresser la mémoire en dehors de l'espace alloué.

Le moniteur a clignoté sauvagement, puis l'écran est devenu noir... jusqu'à ce que Vista démarre le service de réparation du pilote et tout allait bien. Mais il est toujours quelque peu surprenant d'observer si vous avez l'habitude de voir l'erreur typique d'erreur de segmentation sur des programmes standard comme le nôtre. Enfin, une petite critique envers nVidia : dans toute la documentation disponible pour CUDA, il n'y a pas de petit guide qui vous guiderait étape par étape sur la façon de configurer un environnement de développement pour Visual Studio. En fait, le problème n'est pas grave, car le SDK contient un ensemble complet d'exemples que vous pouvez étudier pour comprendre le cadre des applications CUDA, mais un guide du débutant serait utile.

Cliquez sur l'image pour agrandir.

Nvidia a présenté CUDA avec la GeForce 8800. Alors que les promesses étaient tentantes à l'époque, nous avons maintenu notre enthousiasme jusqu'au test proprement dit. En effet, à l'époque, il apparaissait plutôt comme un balisage territorial de rester sur la vague GPGPU. Sans un SDK disponible, il est difficile de dire qu'il ne s'agit pas simplement d'un autre mannequin marketing qui échouera. Ce n'est pas la première fois qu'une bonne initiative est annoncée trop tôt et n'a pas vu le jour à l'époque en raison d'un manque de soutien - surtout dans un secteur aussi concurrentiel. Aujourd'hui, un an et demi après l'annonce, nous pouvons affirmer avec confiance que nVidia a tenu parole.

Le SDK est apparu assez rapidement en bêta début 2007 et a été rapidement mis à jour depuis, ce qui prouve l'intérêt de ce projet pour nVidia. Aujourd'hui CUDA évolue très bien : le SDK est déjà disponible en bêta 2.0 pour les principaux systèmes d'exploitation (Windows XP et Vista, Linux et 1.1 pour Mac OS X), et nVidia a dédié toute une section du site aux développeurs.

Sur un plan plus professionnel, l'impression des premiers pas avec CUDA a été très positive. Même si vous êtes familier avec l'architecture GPU, vous pouvez facilement la comprendre. Lorsque l'API semble simple à première vue, vous commencez immédiatement à croire que vous obtiendrez des résultats convaincants. Mais le temps de calcul ne sera-t-il pas gâché par les nombreux transferts du CPU vers le GPU ? Et comment utiliser ces milliers de threads avec peu ou pas de primitive de synchronisation ? Nous avons commencé nos expériences avec toutes ces peurs en tête. Mais ils se sont rapidement dissipés lorsque la première version de notre algorithme, bien que très triviale, s'est avérée nettement plus rapide que sur le CPU.

Donc CUDA n'est pas une bouée de sauvetage pour les chercheurs qui veulent convaincre les responsables universitaires de leur acheter une GeForce. CUDA est une technologie déjà entièrement disponible que tout programmeur C peut utiliser s'il est prêt à consacrer du temps et des efforts à s'habituer au nouveau paradigme de programmation. Cet effort ne sera pas vain si vos algorithmes se parallélisent bien. Nous tenons également à remercier nVidia pour avoir fourni une documentation complète et de haute qualité aux aspirants programmeurs CUDA pour trouver des réponses.

De quoi CUDA a-t-il besoin pour devenir une API reconnaissable ? En un mot : portabilité. Nous savons que l'avenir de l'informatique réside dans l'informatique parallèle - aujourd'hui, tout le monde se prépare à de tels changements, et toutes les initiatives, tant logicielles que matérielles, sont orientées dans cette direction. Cependant, à l'heure actuelle, si l'on regarde l'évolution des paradigmes, on en est encore au stade initial : on crée des flux manuellement et on essaie de planifier l'accès aux ressources partagées ; vous pouvez en quelque sorte faire face à tout cela, si le nombre de noyaux peut être compté sur les doigts d'une main. Mais dans quelques années, quand le nombre de processeurs se comptera par centaines, cette possibilité n'existera plus. Avec la sortie de CUDA, nVidia a fait le premier pas vers la résolution de ce problème - mais, bien sûr, cette solution ne convient qu'aux GPU de cette société, et même alors pas à tout le monde. Seuls GF8 et 9 (et leurs dérivés Quadro / Tesla) peuvent fonctionner avec les programmes CUDA aujourd'hui. Et la nouvelle gamme 260/280, bien sûr.

Cliquez sur l'image pour agrandir.

Nvidia peut se vanter d'avoir vendu 70 millions de GPU compatibles CUDA dans le monde, mais ce n'est toujours pas suffisant pour devenir le standard de facto. Tenant compte du fait que les concurrents ne restent pas les bras croisés. AMD propose son propre SDK (Stream Computing), et Intel a annoncé une solution (Ct), bien qu'elle ne soit pas encore disponible. Une guerre des standards s'annonce, et il n'y aura clairement pas de place pour trois concurrents sur le marché jusqu'à ce qu'un autre acteur, comme Microsoft, propose une API partagée, ce qui facilitera bien sûr la vie des développeurs.

Par conséquent, nVidia a beaucoup de difficultés à obtenir l'approbation CUDA. Si technologiquement nous avons, sans aucun doute, une solution réussie, il reste encore à convaincre les développeurs de ses perspectives - et ce ne sera pas facile. Cependant, à en juger par les nombreuses annonces et actualités récentes de l'API, l'avenir semble loin d'être sombre.

Selon la théorie de l'évolution de Darwin, les premiers grands singes (si
pour être précis - homo antecessor, prédécesseur humain) est devenu plus tard
en nous. Centres de calcul multitonnes avec un millier de tubes radio ou plus,
occupant des pièces entières, ont été remplacés par des ordinateurs portables d'un demi-kilogramme, qui, soit dit en passant,
ne cédera pas au premier en performance. Les machines à écrire antédiluviennes sont devenues
en imprimant n'importe quoi et n'importe quoi (même sur le corps humain)
appareils multifonctions. Les géants des processeurs ont soudainement décidé de se murer
noyau graphique en "stone". Et les cartes vidéo ont commencé non seulement à montrer une image avec
qualité FPS et graphique acceptable, mais également effectuer toutes sortes de calculs. Oui
comment produire ! A propos de la technologie de calcul multithread au moyen de GPU, et sera discuté.

Pourquoi GPU ?

Je me demande pourquoi ils ont décidé de déplacer toute la puissance de calcul vers le graphique
adaptateur? Comme vous pouvez le constater, les processeurs sont toujours en vogue et il est peu probable qu'ils abandonnent leur chaleur
endroit. Mais le GPU a quelques atouts dans sa manche, avec le joker et les manches.
assez. Le processeur central moderne est affûté pour obtenir le maximum
performances lors du traitement de données entières et de données flottantes
virgule, sans se soucier du traitement parallèle de l'information. Au même
temps, l'architecture de la carte vidéo permet de "paralléliser" rapidement et facilement
traitement de l'information. D'une part, les polygones sont comptés (grâce au convoyeur 3D),
d'autre part, le traitement de la texture des pixels. On peut voir qu'il y a un « bien coordonné
panne "charge dans le noyau de la carte. De plus, la mémoire et le processeur vidéo fonctionnent
plus optimale que la combinaison "RAM-cache-processor". Le moment où l'unité de données
dans la carte vidéo commence à être traité par un processeur de flux GPU, un autre
une unité est chargée en parallèle dans une autre, et, en principe, il est facile à réaliser
Charge GPU comparable à la bande passante du bus,
cependant, pour cela, le chargement des convoyeurs doit être effectué de manière uniforme, sans
tout branchement conditionnel et branchement. Le processeur central, de par sa
l'universalité nécessite une cache pleine de
informations.

Les experts ont réfléchi au travail des GPU dans le calcul parallèle et
mathématiques et déduit la théorie selon laquelle de nombreux calculs scientifiques sont à bien des égards similaires à
traitement de graphiques 3D. De nombreux experts pensent qu'un facteur fondamental de
développement GPGPU (Calcul à usage général sur GPU - universel
calculs au moyen d'une carte vidéo) a été l'émergence en 2003 du projet Brook GPU.

Les créateurs du projet de l'université de Stanford ont dû résoudre un problème
problème : le matériel et le logiciel forcent la carte graphique à produire
calculs polyvalents. Et ils l'ont fait. En utilisant le langage C générique,
Des scientifiques américains ont forcé le GPU à fonctionner comme un processeur, ajusté pour
traitement parallèle. Après Brook, un certain nombre de projets VGA sont apparus,
tels que la bibliothèque Accelerator, la bibliothèque Brahma, le système
métaprogrammation GPU++ et autres.

CUDA !

Le pressentiment des perspectives de développement fait DMLA et NVIDIA
s'accrocher au GPU Brook comme un pit-bull. Si nous omettons la politique de marketing, alors,
après avoir tout mis en œuvre correctement, vous pouvez prendre pied non seulement dans le secteur graphique
marché, mais aussi en informatique (regardez les cartes informatiques spéciales et
les serveurs Tesla avec des centaines de multiprocesseurs), remplaçant les processeurs habituels.

Naturellement, les « maîtres FPS » se sont dispersés devant la pierre d'achoppement, chacun à leur manière.
chemin, mais le principe de base est resté inchangé - effectuer des calculs
Outils GPU. Et maintenant, nous allons examiner de plus près la technologie du "vert" - CUDA
(Architecture de périphérique unifiée de calcul).

Le travail de notre "héroïne" est de fournir une API, et deux à la fois.
Le premier est le runtime CUDA de haut niveau, qui sont des fonctions qui
sont décomposés en niveaux plus simples et transmis à l'API inférieure - pilote CUDA. Donc
que l'expression « haut niveau » s'applique à un processus d'étirement. Tout le sel est
il est dans le pilote, et les bibliothèques, gentiment créées par
développeurs NVIDIA: CUBLAS (outils de calculs mathématiques) et
FFT (calcul à l'aide de l'algorithme de Fourier). Bon, passons à la pratique
parties du matériau.

Terminologie CUDA

NVIDIA fonctionne avec des définitions très particulières pour l'API CUDA. Ils
diffèrent des définitions utilisées pour travailler avec le processeur central.

Fil- le jeu de données à traiter (non
nécessite beaucoup de ressources de traitement).

Chaîne- un groupe de 32 flux. Les données sont traitées uniquement
warps, donc warp est la quantité minimale de données.

Bloquer- un ensemble de flux (de 64 à 512) ou un ensemble
chaînes (de 2 à 16).

Grille Est une collection de blocs. Cette séparation des données
est utilisé uniquement pour améliorer les performances. Donc, si le nombre
multiprocesseurs est volumineux, alors les blocs seront exécutés en parallèle. Si avec
la carte n'a pas eu de chance (les développeurs recommandent d'utiliser
adaptateur n'est pas inférieur au niveau de GeForce 8800 GTS 320 Mo), alors les blocs de données seront traités
régulièrement.

NVIDIA introduit également des concepts tels que noyau, hôte
et dispositif.

Nous travaillons!

Pour travailler pleinement avec CUDA, vous avez besoin de :

1. Connaître la structure des cœurs de shader GPU, comme l'essence de la programmation
est de répartir uniformément la charge entre eux.
2. Être capable de programmer dans l'environnement C, en tenant compte de certains aspects.

Développeurs NVIDIA a révélé les "intérieurs" de la carte vidéo plusieurs
différemment de ce que nous avons l'habitude de voir. Alors bon gré mal gré faut tout étudier
subtilités de l'architecture. Analysons la structure de la "pierre" G80 du légendaire GeForce 8800
GTX.

Le shader core se compose de huit TPC (Texture Processor Cluster) - clusters
processeurs de texture (par exemple, GeForce GTX 280- 15 noyaux, 8800 GTS
il y en a six, 8600 - quatre, etc.). Ceux-ci, à leur tour, se composent de deux
multiprocesseurs de streaming (ci-après dénommés SM). SM (il y a
16) se compose d'un frontal (résout le problème de lecture et de décodage des instructions) et
pipelines back-end (sortie finale des instructions), ainsi que huit SP scalaires (shader
processeur) et deux SFU (unités super fonctionnelles). Pour chaque mesure (unité
time) le frontal sélectionne la chaîne et la traite. À tous les ruisseaux de la chaîne
(rappelez-vous, il y en a 32) ont été traités, 32/8 = 4 cycles sont nécessaires à la fin du pipeline.

Chaque multiprocesseur possède une mémoire dite partagée.
Sa taille est de 16 kilo-octets et offre au programmeur une liberté totale.
action. Distribuez comme vous voulez :). La mémoire partagée fournit une communication de thread dans
un bloc et n'est pas conçu pour fonctionner avec les pixel shaders.

Les SM peuvent également accéder à la GDDR. Pour ce faire, ils ont été "cousus" sur 8 kilo-octets
mémoire cache, stockant tous les éléments les plus importants pour le travail (par exemple, l'informatique
constantes).

Le multiprocesseur a 8192 registres. Le nombre de blocs actifs ne peut pas être
plus de huit, et le nombre de chaînes n'est pas supérieur à 768/32 = 24. On peut voir à partir de là que G80
peut traiter un maximum de 32 * 16 * 24 = 12288 threads par unité de temps. Il est impossible non
tenir compte de ces chiffres lors de l'optimisation du programme à l'avenir (sur un plateau
- la taille du bloc, d'autre part - le nombre de threads). L'équilibre des paramètres peut jouer
rôle important à l'avenir, donc NVIDIA recommande d'utiliser des blocs
avec 128 ou 256 fils. Un bloc de 512 threads est inefficace, car il a
des retards accrus. Compte tenu de toutes les subtilités de la structure de la carte vidéo GPU plus
bonnes compétences en programmation, vous pouvez créer un environnement très productif
outil de calcul parallèle. En parlant de programmation...

La programmation

Pour la "créativité" avec CUDA, il est nécessaire La carte vidéo GeForce n'est pas inférieure
huitième série... AVEC

le site officiel, vous devez télécharger trois packages logiciels : un pilote avec
Prise en charge de CUDA (pour chaque système d'exploitation - le sien), directement le package SDK CUDA (le deuxième
bêta) et des bibliothèques supplémentaires (boîte à outils CUDA). Supports technologiques
systèmes d'exploitation Windows (XP et Vista), Linux et Mac OS X. Pour étudier, je
choisi Vista Ultimate Edition x64 (pour l'avenir, je dirai que le système s'est comporté
Juste parfait). Au moment d'écrire ces lignes, les éléments suivants étaient pertinents pour le travail
Pilote ForceWare 177.35. La boîte à outils utilisée était
Progiciel Borland C ++ 6 Builder (bien que tout environnement fonctionnant avec
langue C).

Il sera facile pour une personne qui connaît la langue de s'habituer au nouvel environnement. Il faut seulement
rappelez-vous les paramètres de base. Le mot-clé _global_ (placé avant la fonction)
indique que la fonction appartient au noyau. Il sera appelé par la centrale
processeur, et tout le travail se fera sur le GPU. L'appel à _global_ nécessite plus
des détails spécifiques, à savoir la taille des mailles, la taille des blocs et quel noyau sera
appliqué. Par exemple, la ligne _global_ void saxpy_parallel<<>>, où X -
la taille de la grille, et Y est la taille du bloc, définit ces paramètres.

Le symbole _device_ signifie que la fonction sera appelée par le cœur graphique, il est
suivra toutes les instructions. Cette fonction est située dans la mémoire du multiprocesseur,
par conséquent, il est impossible d'obtenir son adresse. Le préfixe _host_ signifie que l'appel
et le traitement n'aura lieu qu'avec la participation de la CPU. Il convient de garder à l'esprit que _global_ et
_device_ ne peut pas s'appeler et ne peut pas s'appeler.

De plus, le langage pour CUDA a un certain nombre de fonctions pour travailler avec la mémoire vidéo : cudafree
(libération de mémoire entre GDDR et RAM), cudamemcpy et cudamemcpy2D (copie
mémoire entre GDDR et RAM) et cudamalloc (allocation de mémoire).

Tous les codes de programme sont compilés par l'API CUDA. Le premier est pris
code qui est exclusivement pour le CPU et est exposé
compilation standard et autre code pour l'adaptateur graphique,
est réécrit dans le langage intermédiaire PTX (très similaire à l'assembleur) pour
identifier les erreurs possibles. Après toutes ces "danses", la finale
traduction (traduction) des commandes dans un langage compréhensible pour le GPU/CPU.

Ensemble d'étude

Presque tous les aspects de la programmation sont décrits dans la documentation qui va
accompagné d'un pilote et de deux applications, ainsi que sur le site développeur. Taille
l'article ne suffit pas à les décrire (le lecteur intéressé devra joindre
un peu d'effort et étudiez le matériel par vous-même).

Le navigateur CUDA SDK est spécialement conçu pour les débutants. N'importe qui peut
ressentez la puissance du calcul parallèle sur votre propre peau (le meilleur test pour
stabilité - travail d'exemples sans artefacts et plantages). L'application a
un grand nombre de mini-programmes démonstratifs (61 "tests"). Chaque expérience a
documentation détaillée du code et PDF. Il est immédiatement clair que les gens
ceux qui présentent leurs créations dans le navigateur font un travail sérieux.
Vous pouvez également comparer la vitesse du processeur et de la carte vidéo pendant le traitement
Les données. Par exemple, numériser des tableaux multidimensionnels avec une carte vidéo GeForce 8800
GT 512 Mo avec un bloc de 256 threads produit en 0,17109 millisecondes.
La technologie ne reconnaît pas les tandems SLI, donc si vous avez un duo ou un trio,
désactiver la fonction "couplage" avant le travail, sinon CUDA n'en verra qu'un
dispositif. Double noyau AMD Athlon 64X2(fréquence de base 3000 MHz) même expérience
passe en 2,761528 millisecondes. Il s'avère que le G92 est plus de 16 fois
plus rapide que "pierre" DMLA! Comme vous pouvez le voir, c'est loin d'être un système extrême dans
en tandem avec le système d'exploitation mal aimé des masses, il montre bien
résultats.

En plus du navigateur, il existe un certain nombre de programmes utiles à la société. Adobe
a adapté ses produits aux nouvelles technologies. Photoshop CS4 est maintenant terminé
utilise au moins les ressources des adaptateurs graphiques (vous devez télécharger un
brancher). Des programmes tels que Badaboom media converter et RapiHD peuvent
décoder la vidéo au format MPEG-2. Pas mal pour le traitement du son
utilitaire gratuit Accelero fera l'affaire. Le nombre de logiciels adaptés à l'API CUDA,
grandira sans aucun doute.

Et à ce moment là...

En attendant, vous lisez ce matériel, les travailleurs acharnés des problèmes de processeur
développent leurs propres technologies pour intégrer les GPU dans les CPU. Du côté DMLA tous
compréhensible : ils ont une formidable expérience acquise avec ATI.

La création des « micro-devises », Fusion, sera constituée de plusieurs cœurs sous
nom de code Bulldozer et puce vidéo RV710 (Kong). Leur relation sera
réalisée grâce au bus HyperTransport amélioré. En fonction de la
le nombre de cœurs et leurs caractéristiques de fréquence AMD envisage de créer un prix entier
hiérarchie des "pierres". Il est également prévu de produire des processeurs pour ordinateurs portables (Falcon),
et pour les gadgets multimédias (Bobcat). De plus, c'est l'utilisation de la technologie
dans les appareils portables sera le premier défi pour les Canadiens. Avec le développement
calcul parallèle, l'utilisation de telles "pierres" devrait être très populaire.

Intelligence légèrement en retard avec son Larrabee. Des produits DMLA,
si rien ne se passe, apparaîtra sur les tablettes des magasins fin 2009 - début
2010 année. Et la décision de l'ennemi n'apparaîtra qu'après presque deux
de l'année.

Larrabee aura un grand nombre (lire des centaines) de cœurs. initialement
il y aura des produits conçus pour 8 à 64 cœurs. Ils sont très similaires au Pentium, mais
à peu près retravaillé. Chaque cœur a 256 kilo-octets de cache L2
(il augmentera de taille avec le temps). La relation se fera à travers
Bus en anneau bidirectionnel 1024 bits. Intel dit que leur "enfant"
fonctionne très bien avec DirectX et Open GL API (pour Yabloko), donc non
aucune intervention logicielle n'est requise.

Pourquoi je t'ai dit tout ça ? Evidemment Larrabee et Fusion ne supplanteront pas
processeurs conventionnels et stationnaires du marché, tout comme ils ne seront pas expulsés du marché
cartes vidéo. Pour les gamers et les amoureux de l'extrême, le rêve ultime restera
CPU multicœur et un tandem de plusieurs VGA haut de gamme. Mais quoi encore
les entreprises de traitement passent à l'informatique parallèle basée sur des principes
similaire à GPGPU, en dit long. En particulier, qu'un tel
une technologie comme CUDA a le droit d'exister et est susceptible d'être
très populaire.

Petit résumé

L'informatique parallèle avec une carte graphique n'est qu'un bon outil
entre les mains d'un programmeur assidu. Processeurs improbables dirigés par la loi de Moore
la fin viendra. Entreprises NVIDIA il y a encore un long chemin à parcourir
promouvoir son API auprès des masses (on peut en dire autant de l'idée originale ATI/AMD).
Ce qu'il sera, l'avenir le montrera. Alors CUDA sera de retour :).

P.S. Je recommande aux programmeurs débutants et aux personnes intéressées de visiter
les « établissements virtuels » suivants :

Site Web et site Web officiels de NVIDIA
GPGPU.com. Tous
les informations fournies sont en anglais, mais merci au moins que vous n'êtes pas sur
Chinois. Alors allez-y ! J'espère que l'auteur vous a aidé au moins un peu à
des efforts passionnants pour apprendre CUDA !