Les dispositifs permettant de transformer des ordinateurs personnels en petits superordinateurs sont connus depuis un certain temps. Dans les années 80 du siècle dernier, des transputeurs étaient proposés sur le marché, qui étaient insérés dans les connecteurs d'extension ISA alors courants. Au début, leurs performances dans les tâches pertinentes étaient impressionnantes, mais ensuite la croissance de la vitesse des processeurs universels s'est accélérée, ils ont renforcé leur position dans le calcul parallèle et les transordinateurs n'avaient aucun sens. Bien que des dispositifs similaires existent encore aujourd'hui, il s'agit d'une variété d'accélérateurs spécialisés. Mais souvent, leur champ d’application est restreint et ces accélérateurs ne sont pas particulièrement répandus.

Mais récemment, le relais de l'informatique parallèle est passé au marché de masse, d'une manière ou d'une autre liée aux jeux 3D. Les appareils à usage général dotés de processeurs multicœurs pour les calculs vectoriels parallèles utilisés dans les graphiques 3D atteignent des performances de pointe élevées que les processeurs à usage général ne peuvent égaler. Bien entendu, la vitesse maximale n'est atteinte que dans un certain nombre de tâches pratiques et présente certaines limites, mais de tels dispositifs ont déjà commencé à être utilisés assez largement dans des domaines pour lesquels ils n'étaient pas initialement destinés. Un excellent exemple d'un tel processeur parallèle est le processeur Cell, développé par l'alliance Sony-Toshiba-IBM et utilisé dans la console de jeu Sony PlayStation 3, ainsi que dans toutes les cartes vidéo modernes des leaders du marché - Nvidia et AMD.

Nous ne toucherons pas à Cell aujourd'hui, bien qu'il soit apparu plus tôt et qu'il s'agisse d'un processeur universel doté de capacités vectorielles supplémentaires, nous n'en parlons pas aujourd'hui. Pour les accélérateurs vidéo 3D, sont apparues il y a quelques années les premières technologies de calculs généraux non graphiques GPGPU (General-Purpose computing on GPUs). Après tout, les puces vidéo modernes contiennent des centaines d’unités d’exécution mathématique, et cette puissance peut être utilisée pour accélérer considérablement de nombreuses applications gourmandes en calcul. Et les générations actuelles de GPU ont une architecture assez flexible qui, associée à des langages de programmation de haut niveau et à des architectures matérielles-logicielles comme celle évoquée dans cet article, révèle ces capacités et les rend beaucoup plus accessibles.

La création de GPCPU a été motivée par l'émergence de programmes de shader assez rapides et flexibles pouvant être exécutés par des puces vidéo modernes. Les développeurs ont décidé de faire en sorte que les GPU calculent non seulement des images dans des applications 3D, mais également qu'ils soient utilisés dans d'autres calculs parallèles. Dans GPGPU, des API graphiques ont été utilisées pour cela : OpenGL et Direct3D, lorsque les données étaient transférées vers la puce vidéo sous forme de textures, et que les programmes de calcul étaient chargés sous forme de shaders. Les inconvénients de cette méthode sont la complexité de programmation relativement élevée, la faible vitesse d'échange de données entre le CPU et le GPU et d'autres limitations, dont nous parlerons plus tard.

Le calcul GPU s’est développé et se développe très rapidement. Par la suite, deux grands fabricants de puces vidéo, Nvidia et AMD, ont développé et annoncé des plates-formes correspondantes appelées respectivement CUDA (Compute Unified Device Architecture) et CTM (Close To Metal ou AMD Stream Computing). Contrairement aux modèles de programmation GPU précédents, ceux-ci ont été conçus avec un accès direct aux capacités matérielles des cartes vidéo. Les plates-formes ne sont pas compatibles entre elles, CUDA est une extension du langage de programmation C et CTM est une machine virtuelle qui exécute du code assembleur. Mais les deux plates-formes ont éliminé certaines des limitations importantes des modèles GPGPU précédents qui utilisaient un pipeline graphique traditionnel et les interfaces Direct3D ou OpenGL correspondantes.

Bien entendu, les standards ouverts utilisant OpenGL semblent être les plus portables et les plus universels, permettant d’utiliser le même code sur des puces vidéo de différents fabricants. Mais ces méthodes présentent de nombreux inconvénients : elles sont beaucoup moins flexibles et moins pratiques à utiliser. De plus, ils ne permettent pas d'utiliser les capacités spécifiques de certaines cartes vidéo, telles que la mémoire partagée (commune) rapide présente dans les processeurs informatiques modernes.

C'est pourquoi Nvidia a lancé la plateforme CUDA, un langage de programmation de type C avec son propre compilateur et ses propres bibliothèques pour le calcul GPU. Bien sûr, écrire du code optimal pour les puces vidéo n'est pas du tout si simple et cette tâche nécessite beaucoup de travail manuel, mais CUDA révèle toutes les possibilités et donne au programmeur plus de contrôle sur les capacités matérielles du GPU. Il est important que le support Nvidia CUDA soit disponible dans les puces G8x, G9x et GT2xx utilisées dans les cartes vidéo des séries Geforce 8, 9 et 200, qui sont très répandues. La version finale de CUDA 2.0 est désormais disponible, elle apporte quelques nouvelles fonctionnalités, comme la prise en charge des calculs en double précision. CUDA est disponible sur les systèmes d'exploitation Linux, Windows et MacOS X 32 bits et 64 bits.

Différence entre CPU et GPU dans les calculs parallèles

L'augmentation des fréquences des processeurs universels s'est heurtée à des limitations physiques et à une consommation d'énergie élevée, et leurs performances sont de plus en plus augmentées en plaçant plusieurs cœurs sur une seule puce. Les processeurs actuellement vendus ne contiennent que quatre cœurs (la croissance future ne sera pas rapide) et sont conçus pour des applications générales, utilisant plusieurs commandes et flux de données MIMD. Chaque noyau fonctionne séparément des autres, exécutant différentes instructions pour différents processus.

Des capacités vectorielles spécialisées (SSE2 et SSE3) pour les vecteurs à quatre composants (virgule flottante simple précision) et à deux composants (double précision) sont apparues dans les processeurs à usage général en raison des exigences accrues des applications graphiques, principalement. C'est pourquoi, pour certaines tâches, l'utilisation de GPU est plus rentable, car ils ont été initialement conçus pour elles.

Par exemple, dans les puces vidéo Nvidia, l'unité principale est un multiprocesseur avec huit à dix cœurs et des centaines d'ALU au total, plusieurs milliers de registres et une petite quantité de mémoire partagée. De plus, la carte vidéo contient une mémoire globale rapide accessible par tous les multiprocesseurs, une mémoire locale dans chaque multiprocesseur et une mémoire spéciale pour les constantes.

Plus important encore, ces multiples cœurs multiprocesseurs du GPU sont des cœurs SIMD (single instruction stream, multiple data stream). Et ces cœurs exécutent simultanément les mêmes instructions, un style de programmation commun aux algorithmes graphiques et à de nombreuses tâches scientifiques, mais qui nécessite une programmation spécifique. Mais cette approche permet d'augmenter le nombre d'unités d'exécution grâce à leur simplification.

Listons donc les principales différences entre les architectures CPU et GPU. Les cœurs de processeur sont conçus pour exécuter un seul flux d'instructions séquentielles avec des performances maximales, tandis que les GPU sont conçus pour exécuter rapidement un grand nombre de flux d'instructions parallèles. Les processeurs à usage général sont optimisés pour obtenir des performances élevées à partir d'un seul flux d'instructions, traitant à la fois des nombres entiers et des nombres à virgule flottante. Dans ce cas, l'accès à la mémoire est aléatoire.

Les développeurs de processeurs essaient d'exécuter autant d'instructions que possible en parallèle pour augmenter les performances. Pour y parvenir, à commencer par les processeurs Intel Pentium, une exécution superscalaire est apparue, assurant l'exécution de deux instructions par cycle d'horloge, et le Pentium Pro se distinguait par une exécution des instructions dans le désordre. Mais l'exécution parallèle d'un flux séquentiel d'instructions présente certaines limites fondamentales et l'augmentation du nombre d'unités d'exécution ne peut pas entraîner une augmentation multiple de la vitesse.

Les puces vidéo ont dès le départ un fonctionnement simple et parallèle. La puce vidéo prend un groupe de polygones en entrée, effectue toutes les opérations nécessaires et produit des pixels en sortie. Le traitement des polygones et des pixels est indépendant ; ils peuvent être traités en parallèle, séparément les uns des autres. Par conséquent, en raison de l'organisation intrinsèquement parallèle du travail, le GPU utilise un grand nombre d'unités d'exécution, faciles à charger, contrairement au flux séquentiel d'instructions pour le CPU. De plus, les GPU modernes peuvent également exécuter plus d’une instruction par cycle d’horloge (double problème). Ainsi, l’architecture Tesla, sous certaines conditions, lance simultanément les opérations MAD+MUL ou MAD+SFU.

Le GPU diffère également du CPU en termes de principes d'accès à la mémoire. Sur le GPU, cela est cohérent et facilement prévisible - si un texel de texture est lu dans la mémoire, après un certain temps, le temps viendra pour les texels voisins. Et lors de l'enregistrement, la même chose se produit : un pixel est écrit dans le framebuffer, et après quelques cycles d'horloge, celui situé à côté sera enregistré. L’organisation de la mémoire est donc différente de celle utilisée dans le CPU. Et la puce vidéo, contrairement aux processeurs universels, n'a tout simplement pas besoin d'une grande mémoire cache, et les textures ne nécessitent que quelques kilo-octets (jusqu'à 128-256 dans les GPU actuels).

Et le travail avec la mémoire elle-même est quelque peu différent pour les GPU et les CPU. Ainsi, tous les processeurs centraux n'ont pas de contrôleurs de mémoire intégrés, et tous les GPU ont généralement plusieurs contrôleurs, jusqu'à huit canaux 64 bits dans la puce Nvidia GT200. De plus, les cartes vidéo utilisent une mémoire plus rapide et, par conséquent, les puces vidéo ont accès à une bande passante mémoire nettement plus importante, ce qui est également très important pour les calculs parallèles fonctionnant avec d'énormes flux de données.

Dans les processeurs universels, un grand nombre de transistors et de zones de puce sont utilisés pour les tampons d'instructions, la prédiction de branchement matériel et d'énormes quantités de mémoire cache sur puce. Tous ces blocs matériels sont nécessaires pour accélérer l'exécution de quelques flux de commandes. Les puces vidéo utilisent des transistors pour des réseaux d'unités d'exécution, des unités de contrôle de flux, une petite mémoire partagée et des contrôleurs de mémoire pour plusieurs canaux. Ce qui précède n'accélère pas l'exécution des threads individuels, mais permet à la puce de gérer plusieurs milliers de threads exécutés simultanément par la puce et nécessitant une bande passante mémoire élevée.

À propos des différences de mise en cache. Les processeurs à usage général utilisent la mémoire cache pour augmenter les performances en réduisant la latence d'accès à la mémoire, tandis que les GPU utilisent le cache ou la mémoire partagée pour augmenter la bande passante. Les processeurs réduisent la latence d'accès à la mémoire en utilisant des caches volumineux et la prédiction de branche de code. Ces éléments matériels occupent la majeure partie de la surface de la puce et consomment beaucoup d’énergie. Les puces vidéo contournent le problème des délais d'accès à la mémoire en exécutant des milliers de threads simultanément - pendant que l'un des threads attend des données de la mémoire, la puce vidéo peut effectuer des calculs par un autre thread sans attente ni retard.

Il existe également de nombreuses différences dans la prise en charge du multithreading. Le processeur exécute 1 à 2 threads de calcul par cœur de processeur, et les puces vidéo peuvent prendre en charge jusqu'à 1 024 threads par multiprocesseur, dont plusieurs sur la puce. Et si le passage d'un thread à un autre coûte des centaines de cycles d'horloge au processeur, alors le GPU commute plusieurs threads en un seul cycle d'horloge.

De plus, les processeurs utilisent des unités SIMD (single instruction over multiple data) pour les calculs vectoriels, et les GPU utilisent SIMT (single instruction multiple threads) pour le traitement des threads scalaires. SIMT n'oblige pas le développeur à convertir les données en vecteurs et permet un branchement arbitraire dans les flux.

Bref, on peut dire que contrairement aux processeurs universels modernes, les puces vidéo sont conçues pour des calculs parallèles avec un grand nombre d'opérations arithmétiques. Et un nombre beaucoup plus important de transistors GPU fonctionnent conformément à leur destination - traiter des tableaux de données, et ne contrôlent pas l'exécution (contrôle de flux) de quelques threads de calcul séquentiels. Voici un diagramme de l'espace occupé par diverses logiques sur le CPU et le GPU :

En conséquence, la base pour utiliser efficacement la puissance du GPU dans les calculs scientifiques et autres calculs non graphiques est la parallélisation des algorithmes sur des centaines d’unités d’exécution disponibles sur les puces vidéo. Par exemple, de nombreuses applications de modélisation moléculaire sont bien adaptées aux calculs sur puces vidéo ; elles nécessitent une grande puissance de calcul et sont donc pratiques pour le calcul parallèle. Et l’utilisation de plusieurs GPU fournit encore plus de puissance de calcul pour résoudre de tels problèmes.

L'exécution de calculs sur le GPU montre d'excellents résultats dans les algorithmes utilisant le traitement parallèle des données. C’est-à-dire lorsque la même séquence d’opérations mathématiques est appliquée à une grande quantité de données. Dans ce cas, les meilleurs résultats sont obtenus si le rapport entre le nombre d'instructions arithmétiques et le nombre d'accès mémoire est suffisamment grand. Cela impose moins d'exigences en matière de contrôle de flux, et la haute densité mathématique et le grand volume de données éliminent le besoin de caches volumineux, comme pour le processeur.

En raison de toutes les différences décrites ci-dessus, les performances théoriques des puces vidéo dépassent largement celles du processeur. Nvidia fournit le graphique suivant de la croissance des performances du CPU et du GPU au cours des dernières années :

Naturellement, ces données ne sont pas sans part de ruse. Après tout, sur un CPU, il est en pratique beaucoup plus facile d'obtenir des chiffres théoriques, et les chiffres sont donnés en simple précision dans le cas d'un GPU, et en double précision dans le cas d'un CPU. Dans tous les cas, pour certaines tâches parallèles, une simple précision suffit, et la différence de vitesse entre les processeurs universels et graphiques est très grande, et le jeu en vaut donc la chandelle.

Premières tentatives d'utilisation des calculs GPU

Cela fait un certain temps qu'ils essaient d'utiliser des puces vidéo dans des calculs mathématiques parallèles. Les toutes premières tentatives de telles applications étaient extrêmement primitives et se limitaient à l'utilisation de certaines fonctions matérielles, telles que la rastérisation et le Z-buffering. Mais au cours du siècle actuel, avec l’avènement des shaders, les calculs matriciels ont commencé à s’accélérer. En 2003, chez SIGGRAPH, une section distincte a été réservée au calcul GPU, et elle s'appelait GPGPU (General-Purpose computing on GPU).

Le plus connu est BrookGPU, le compilateur du langage de programmation Brook Stream, conçu pour effectuer des calculs non graphiques sur le GPU. Avant son apparition, les développeurs utilisant les capacités des puces vidéo pour les calculs choisissaient l'une des deux API courantes : Direct3D ou OpenGL. Cela a sérieusement limité l'utilisation des GPU, car les graphiques 3D utilisent des shaders et des textures que les spécialistes de la programmation parallèle ne sont pas tenus de connaître ; ils utilisent des threads et des cœurs. Brook a pu les aider à rendre leur tâche plus facile. Ces extensions de streaming du langage C, développées à l'Université de Stanford, cachaient l'API 3D aux programmeurs et présentaient la puce vidéo comme un coprocesseur parallèle. Le compilateur a traité le fichier .br avec du code et des extensions C++, produisant du code lié à une bibliothèque compatible DirectX, OpenGL ou x86.

Naturellement, Brook présentait de nombreuses lacunes sur lesquelles nous nous sommes attardés et dont nous parlerons plus en détail plus tard. Mais même son apparition a provoqué un regain d'attention significatif de la part des mêmes Nvidia et ATI sur l'initiative du calcul sur GPU, puisque le développement de ces capacités a sérieusement changé le marché à l'avenir, ouvrant un tout nouveau secteur de celui-ci - les ordinateurs parallèles. basé sur des puces vidéo.

Par la suite, certains chercheurs du projet Brook ont ​​rejoint l'équipe de développement de Nvidia pour introduire une stratégie de calcul parallèle matériel-logiciel, ouvrant ainsi de nouvelles parts de marché. Et le principal avantage de cette initiative de Nvidia est que les développeurs connaissent toutes les capacités de leurs GPU dans les moindres détails, et il n'est pas nécessaire d'utiliser l'API graphique, et vous pouvez travailler avec le matériel directement à l'aide du pilote. Le résultat des efforts de cette équipe a été Nvidia CUDA (Compute Unified Device Architecture), une nouvelle architecture logicielle et matérielle pour le calcul parallèle sur GPU Nvidia, qui fait l'objet de cet article.

Domaines d'application des calculs parallèles sur GPU

Pour comprendre l’intérêt du transfert des calculs sur puces vidéo, voici les chiffres moyens obtenus par les chercheurs du monde entier. En moyenne, lors du transfert de calculs vers le GPU, de nombreuses tâches atteignent des vitesses de 5 à 30 fois supérieures à celles des processeurs rapides à usage général. Les plus grands nombres (de l'ordre d'une accélération de 100x voire plus !) sont obtenus avec un code peu adapté aux calculs utilisant des blocs SSE, mais assez pratique pour les GPU.

Ce ne sont là que quelques exemples d'accélérations du code synthétique sur le GPU par rapport au code vectoriel SSE sur le CPU (selon Nvidia) :

• Microscopie à fluorescence : 12x ;
• Dynamique moléculaire (calcul de force non liée) : 8-16x ;
• Électrostatique (sommation coulombienne directe et multiniveau) : 40-120x et 7x.

Et c'est un signe que Nvidia aime beaucoup, le montrant à toutes les présentations, sur lequel nous reviendrons plus en détail dans la deuxième partie de l'article, consacrée à des exemples précis d'applications pratiques de l'informatique CUDA :

Comme vous pouvez le constater, les chiffres sont très attractifs, les multiplications par 100 à 150 sont particulièrement impressionnantes. Dans le prochain article sur CUDA, nous examinerons certains de ces chiffres en détail. Listons maintenant les principales applications dans lesquelles le GPU computing est actuellement utilisé : analyse et traitement d'images et de signaux, simulation physique, mathématiques computationnelles, biologie computationnelle, calculs financiers, bases de données, dynamique des gaz et liquides, cryptographie, radiothérapie adaptative, astronomie, traitement du son, bioinformatique, simulations biologiques, vision par ordinateur, exploration de données, cinéma et télévision numériques, simulations électromagnétiques, systèmes d'information géographique, applications militaires, planification minière, dynamique moléculaire, imagerie par résonance magnétique (IRM), réseaux de neurones, recherche océanographique, physique des particules , simulation de repliement de protéines, chimie quantique, lancer de rayons, visualisation, radar, simulation de réservoir, intelligence artificielle, analyse de données satellitaires, exploration sismique, chirurgie, échographie, vidéoconférence.

Des détails sur de nombreuses applications peuvent être trouvés sur le site Web de Nvidia dans la section. Comme vous pouvez le constater, la liste est assez longue, mais ce n’est pas tout ! Cela peut se poursuivre et on peut certainement supposer que dans le futur d'autres domaines d'application des calculs parallèles sur puces vidéo seront découverts, dont nous n'avons pas encore connaissance.

Fonctionnalités Nvidia CUDA

La technologie CUDA est une architecture informatique matériel-logiciel Nvidia basée sur une extension du langage C, qui permet d'organiser l'accès au jeu d'instructions d'un accélérateur graphique et de gérer sa mémoire lors de l'organisation du calcul parallèle. CUDA aide à mettre en œuvre des algorithmes exécutables sur les accélérateurs vidéo GeForce de huitième génération et plus anciens (Geforce 8, Geforce 9, Geforce 200 series), ainsi que sur Quadro et Tesla.

Bien que la complexité de la programmation des GPU à l’aide de CUDA soit assez élevée, elle est inférieure à celle des solutions GPGPU antérieures. De tels programmes nécessitent de diviser l'application entre plusieurs multiprocesseurs, similaire à la programmation MPI, mais sans partager les données stockées dans la mémoire vidéo partagée. Et comme la programmation CUDA pour chaque multiprocesseur est similaire à la programmation OpenMP, elle nécessite une bonne compréhension de l'organisation de la mémoire. Mais bien entendu, la complexité du développement et du portage vers CUDA dépend grandement de l’application.

Le kit de développement contient de nombreux exemples de code et est bien documenté. Le processus d'apprentissage prendra environ deux à quatre semaines pour ceux qui sont déjà familiarisés avec OpenMP et MPI. L'API est basée sur le langage C étendu, et pour traduire le code de ce langage, le SDK CUDA inclut le compilateur de ligne de commande nvcc, basé sur le compilateur open source Open64.

Listons les principales caractéristiques de CUDA :

• solution logicielle et matérielle unifiée pour le calcul parallèle sur les puces vidéo Nvidia ;
• une large gamme de solutions prises en charge, du mobile au multipuce
• langage de programmation C standard ;
• bibliothèques d'analyse numérique standards FFT (Fast Fourier Transform) et BLAS (Linear Algebra) ;
• échange de données optimisé entre CPU et GPU ;
• interaction avec les API graphiques OpenGL et DirectX ;
• prise en charge des systèmes d'exploitation 32 et 64 bits : Windows XP, Windows Vista, Linux et MacOS X ;
• Possibilité d'aménagement à bas niveau.

Concernant la prise en charge du système d'exploitation, il convient d'ajouter que toutes les principales distributions Linux sont officiellement prises en charge (Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), mais, à en juger par les données des passionnés, CUDA fonctionne très bien sur d'autres versions : Fedora Core, Ubuntu, Gentoo, etc.

L'environnement de développement CUDA (CUDA Toolkit) comprend :

• compilateur nvcc ;
• Bibliothèques FFT et BLAS ;
• profileur ;
• débogueur gdb pour GPU ;
• Pilote d'exécution CUDA inclus avec les pilotes Nvidia standard
• manuel de programmation ;
• SDK développeur CUDA (code source, utilitaires et documentation).

Dans le code source exemples : tri bitonique parallèle, transposition matricielle, sommation de préfixes parallèles de grands tableaux, convolution d'image, transformée en ondelettes discrète, exemple d'interaction avec OpenGL et Direct3D, utilisation des bibliothèques CUBLAS et CUFFT, calcul du prix de l'option (Black formule Scholes, modèle binomial, méthode Monte Carlo), générateur de nombres aléatoires parallèles Mersenne Twister, calcul d'histogramme d'un grand tableau, réduction du bruit, filtre Sobel (trouver des limites).

Avantages et limites de CUDA

Du point de vue d'un programmeur, un pipeline graphique est un ensemble d'étapes de traitement. Le bloc de géométrie génère les triangles et le bloc de rastérisation génère les pixels affichés sur le moniteur. Le modèle de programmation GPGPU traditionnel ressemble à ceci :

Pour transférer les calculs vers le GPU au sein de ce modèle, une approche particulière est nécessaire. Même l'ajout élément par élément de deux vecteurs nécessitera de dessiner la figure sur l'écran ou dans un tampon hors écran. La figure est pixellisée, la couleur de chaque pixel est calculée à l'aide d'un programme donné (pixel shader). Le programme lit les données d'entrée des textures pour chaque pixel, les ajoute et les écrit dans le tampon de sortie. Et toutes ces nombreuses opérations sont nécessaires pour quelque chose qui est écrit avec un seul opérateur dans un langage de programmation classique !

Par conséquent, l’utilisation du GPGPU à des fins informatiques générales présente la limite d’être trop difficile à former pour les développeurs. Oui, et il y a suffisamment d'autres restrictions, car un pixel shader n'est qu'une formule pour la dépendance de la couleur finale d'un pixel sur ses coordonnées, et le langage des pixel shaders est un langage pour écrire ces formules avec une syntaxe de type C . Les premières méthodes GPGPU constituent une astuce intéressante qui vous permet d’utiliser la puissance du GPU, mais sans aucune commodité. Les données y sont représentées par des images (textures) et l'algorithme est représenté par le processus de rastérisation. Il convient de noter en particulier le modèle très spécifique de mémoire et d’exécution.

L'architecture logicielle et matérielle de Nvidia pour le calcul GPU diffère des modèles GPGPU précédents en ce qu'elle vous permet d'écrire des programmes pour le GPU dans un vrai langage C avec une syntaxe standard, des pointeurs et la nécessité d'un minimum d'extensions pour accéder aux ressources informatiques des puces vidéo. CUDA est indépendant des API graphiques et possède certaines fonctionnalités conçues spécifiquement pour l'informatique générale.

Avantages de CUDA par rapport à l'approche traditionnelle du calcul GPGPU :

• L'interface de programmation d'applications CUDA est basée sur le langage de programmation standard C avec des extensions, ce qui simplifie le processus d'apprentissage et de mise en œuvre de l'architecture CUDA ;
• CUDA donne accès à 16 Ko de mémoire partagée par thread par multiprocesseur, qui peut être utilisée pour organiser un cache avec une bande passante plus élevée que les récupérations de texture ;
• transfert de données plus efficace entre le système et la mémoire vidéo
• pas besoin d'API graphiques avec redondance et surcharge ;
• adressage, collecte et dispersion de mémoire linéaire, capacité d'écrire à des adresses arbitraires ;
• prise en charge matérielle des opérations sur les entiers et les bits.

Principales limites de CUDA :

• manque de prise en charge de la récursivité pour les fonctions exécutables ;
• largeur de bloc minimale de 32 threads ;
• architecture CUDA fermée appartenant à Nvidia.

Les faiblesses de la programmation avec les méthodes GPGPU précédentes sont que ces méthodes n'utilisent pas d'unités d'exécution de vertex shader dans les architectures non unifiées précédentes, les données sont stockées dans des textures et sorties dans un tampon hors écran, et les algorithmes multi-passes utilisent des unités de pixel shader. Les limitations du GPGPU peuvent inclure : une utilisation insuffisante des capacités matérielles, des limitations de la bande passante mémoire, le manque d'opération de dispersion (collecte uniquement), l'utilisation obligatoire de l'API graphique.

Les principaux avantages de CUDA par rapport aux méthodes GPGPU précédentes proviennent du fait que l'architecture est conçue pour utiliser efficacement le calcul non graphique sur le GPU et utilise le langage de programmation C sans nécessiter le portage des algorithmes vers une forme conviviale pour le concept de pipeline graphique. . CUDA offre une nouvelle voie vers le calcul GPU qui n'utilise pas d'API graphiques, offrant un accès aléatoire à la mémoire (dispersion ou rassemblement). Cette architecture ne présente pas les inconvénients du GPGPU et utilise toutes les unités d'exécution, et étend également les capacités grâce aux mathématiques entières et aux opérations de décalage de bits.

De plus, CUDA ouvre certaines fonctionnalités matérielles non disponibles dans les API graphiques, telles que la mémoire partagée. Il s'agit d'une petite mémoire (16 kilo-octets par multiprocesseur) à laquelle les blocs de threads ont accès. Il vous permet de mettre en cache les données les plus fréquemment consultées et peut fournir des vitesses plus rapides que l'utilisation de récupérations de textures pour cette tâche. Ce qui, à son tour, réduit la sensibilité du débit des algorithmes parallèles dans de nombreuses applications. Par exemple, il est utile pour l'algèbre linéaire, la transformée de Fourier rapide et les filtres de traitement d'image.

L'accès à la mémoire est également plus pratique dans CUDA. Le code de l'API graphique génère des données sous forme de 32 valeurs à virgule flottante simple précision (valeurs RGBA simultanément dans huit cibles de rendu) dans des zones prédéfinies, et CUDA prend en charge l'écriture dispersée - un nombre illimité d'enregistrements à n'importe quelle adresse. De tels avantages permettent d'exécuter certains algorithmes sur le GPU qui ne peuvent pas être implémentés efficacement à l'aide des méthodes GPGPU basées sur des API graphiques.

De plus, les API graphiques stockent nécessairement les données dans des textures, ce qui nécessite un conditionnement préalable de grands tableaux dans des textures, ce qui complique l'algorithme et oblige à utiliser un adressage spécial. Et CUDA vous permet de lire des données à n'importe quelle adresse. Un autre avantage de CUDA est l'échange de données optimisé entre le CPU et le GPU. Et pour les développeurs qui souhaitent un accès de bas niveau (par exemple, lors de l'écriture d'un autre langage de programmation), CUDA offre des capacités de programmation en langage assembleur de bas niveau.

Histoire du développement de CUDA

Le développement de CUDA a été annoncé avec la puce G80 en novembre 2006, et la sortie d'une version bêta publique du SDK CUDA a eu lieu en février 2007. La version 1.0 est sortie en juin 2007 pour lancer les solutions Tesla basées sur la puce G80 et destinées au marché du calcul haute performance. Puis, à la fin de l'année, la version bêta de CUDA 1.1 est sortie, qui, malgré la légère augmentation du numéro de version, a introduit pas mal de nouveautés.

La nouveauté de CUDA 1.1 est l'inclusion de la fonctionnalité CUDA dans les pilotes vidéo Nvidia classiques. Cela signifiait que dans les exigences de tout programme CUDA, il suffisait d'indiquer une carte vidéo de la série Geforce 8 et supérieure, ainsi qu'une version minimale du pilote 169.xx. Ceci est très important pour les développeurs : si ces conditions sont remplies, les programmes CUDA fonctionneront pour n'importe quel utilisateur. L'exécution asynchrone ainsi que la copie de données (uniquement pour les puces G84, G86, G92 et supérieures), le transfert de données asynchrone vers la mémoire vidéo, les opérations d'accès à la mémoire atomique, la prise en charge des versions 64 bits de Windows et la possibilité d'exploiter plusieurs systèmes ont également été ajoutés. puce CUDA en mode SLI.

Pour le moment, la version actuelle est destinée aux solutions basées sur GT200 CUDA 2.0, publiées avec la gamme Geforce GTX 200. La version bêta a été publiée au printemps 2008. La deuxième version ajoutée : prise en charge des calculs double précision (support matériel uniquement pour le GT200), enfin prise en charge de Windows Vista (versions 32 et 64 bits) et Mac OS X, ajout d'outils de débogage et de profilage, prise en charge des textures 3D, transfert de données optimisé .

Quant aux calculs en double précision, leur rapidité sur la génération matérielle actuelle est plusieurs fois inférieure à celle en simple précision. Les raisons sont discutées dans la nôtre. La mise en œuvre de ce support dans le GT200 est que les blocs FP32 ne sont pas utilisés pour obtenir des résultats à une vitesse quatre fois inférieure ; pour prendre en charge les calculs FP64, Nvidia a décidé de créer des unités de calcul dédiées. Et dans le GT200, il y en a dix fois moins que les blocs FP32 (un bloc double précision pour chaque multiprocesseur).

En réalité, les performances peuvent être encore inférieures, puisque l'architecture est optimisée pour la lecture 32 bits de la mémoire et des registres ; de plus, la double précision n'est pas nécessaire dans les applications graphiques, et dans le GT200, elle est plus susceptible d'être simplement là. Et les processeurs quad-core modernes n'affichent pas beaucoup moins de performances réelles. Mais étant même 10 fois plus lent que la simple précision, une telle prise en charge est utile pour les conceptions à précision mixte. Une technique courante consiste à obtenir dans un premier temps des résultats approximatifs en simple précision, puis à les affiner en double précision. Cela peut désormais être fait directement sur la carte vidéo, sans envoyer de données intermédiaires au CPU.

Une autre fonctionnalité utile de CUDA 2.0 n'a rien à voir avec le GPU, curieusement. Simplement, vous pouvez désormais compiler le code CUDA en code SSE multithread très efficace pour une exécution rapide sur le CPU. Autrement dit, cette fonctionnalité convient désormais non seulement au débogage, mais également à une utilisation réelle sur des systèmes sans carte vidéo Nvidia. Après tout, l'utilisation de CUDA dans le code standard est entravée par le fait que les cartes vidéo Nvidia, bien que les plus populaires parmi les solutions vidéo dédiées, ne soient pas disponibles sur tous les systèmes. Et avant la version 2.0, dans de tels cas, il aurait fallu créer deux codes différents : pour CUDA et séparément pour le CPU. Et désormais, vous pouvez exécuter n'importe quel programme CUDA sur un processeur avec une efficacité élevée, bien qu'à une vitesse inférieure à celle des puces vidéo.

Solutions prenant en charge Nvidia CUDA

Toutes les cartes graphiques compatibles CUDA peuvent contribuer à accélérer les tâches les plus exigeantes, du traitement audio et vidéo à la médecine et à la recherche scientifique. La seule véritable limitation est que de nombreux programmes CUDA nécessitent un minimum de 256 Mo de mémoire vidéo, ce qui constitue l'une des caractéristiques techniques les plus importantes des applications CUDA.

La dernière liste des produits prenant en charge CUDA est disponible à l'adresse. Au moment de la rédaction de cet article, les calculs CUDA prenaient en charge tous les produits des séries GeForce 200, GeForce 9 et GeForce 8, y compris les produits mobiles commençant par la GeForce 8400M, ainsi que les chipsets GeForce 8100, 8200 et 8300. Les produits Quadro et 8300 modernes également prend en charge CUDA tous les Tesla : S1070, C1060, C870, D870 et S870.

Notons notamment qu'à côté des nouvelles cartes vidéo Geforce GTX 260 et 280, des solutions correspondantes pour le calcul haute performance ont été annoncées : Tesla C1060 et S1070 (montrées sur la photo ci-dessus), qui seront disponibles à l'achat cet automne. Ils utilisent le même GPU - GT200, dans le C1060 il y en a un, dans le S1070 il y en a quatre. Mais contrairement aux solutions de jeux, elles utilisent quatre gigaoctets de mémoire par puce. Le seul inconvénient est que la fréquence de la mémoire et la bande passante sont inférieures à celles des cartes de jeu, offrant 102 Go/s par puce.

Composition de Nvidia CUDA

CUDA comprend deux API : de haut niveau (API CUDA Runtime) et de bas niveau (API CUDA Driver), bien qu'il soit impossible d'utiliser les deux en même temps dans un seul programme ; vous devez utiliser l'une ou l'autre. Celui de haut niveau fonctionne « au-dessus » de celui de bas niveau, tous les appels d'exécution sont traduits en instructions simples traitées par l'API Driver de bas niveau. Mais même une API « de haut niveau » suppose des connaissances sur la conception et le fonctionnement des puces vidéo Nvidia ; il n’y a pas là de niveau d’abstraction trop élevé.

Il existe un autre niveau, encore plus élevé - deux bibliothèques :

CUBLAS Version CUDA de BLAS (Basic Linear Algebra Subprograms), conçue pour calculer des problèmes d'algèbre linéaire et utiliser un accès direct aux ressources GPU ;

MANCHETTE Version CUDA de la bibliothèque Fast Fourier Transform pour calculer la transformée de Fourier rapide, largement utilisée dans le traitement du signal. Les types de transformation suivants sont pris en charge : complexe-complexe (C2C), réel-complexe (R2C) et complexe-réel (C2R).

Examinons de plus près ces bibliothèques. Les CUBLAS sont des algorithmes d'algèbre linéaire standards traduits en langage CUDA ; actuellement, seul un certain ensemble de fonctions de base de CUBLAS est pris en charge. La bibliothèque est très simple à utiliser : vous devez créer une matrice et des objets vectoriels dans la mémoire de la carte vidéo, les remplir de données, appeler les fonctions CUBLAS requises et charger les résultats de la mémoire vidéo dans la mémoire système. CUBLAS contient des fonctions spéciales pour créer et détruire des objets dans la mémoire GPU, ainsi que pour lire et écrire des données dans cette mémoire. Fonctions BLAS prises en charge : niveaux 1, 2 et 3 pour les nombres réels, CGEMM niveau 1 pour les nombres complexes. Le niveau 1 correspond aux opérations vecteur-vecteur, le niveau 2 correspond aux opérations vecteur-matrice, le niveau 3 correspond aux opérations matrice-matrice.

CUFFT Version CUDA de la fonction Fast Fourier Transform, largement utilisée et très importante dans l'analyse du signal, le filtrage, etc. CUFFT fournit une interface simple pour calculer efficacement la FFT sur les GPU Nvidia sans avoir besoin de développer votre propre GPU FFT. La variante CUDA de FFT prend en charge les transformations 1D, 2D et 3D de données complexes et réelles, l'exécution par lots pour plusieurs transformations 1D en parallèle, les tailles des transformations 2D et 3D peuvent être inférieures à , pour 1D, la taille allant jusqu'à 8 millions d'éléments est prise en charge.

Bases de la création de programmes sur CUDA

Pour comprendre la suite du texte, vous devez comprendre les caractéristiques architecturales de base des puces vidéo Nvidia. Le GPU est constitué de plusieurs clusters d'unités de texture (Texture Processing Cluster). Chaque cluster se compose d'un grand bloc de récupérations de textures et de deux ou trois multiprocesseurs de streaming, chacun étant composé de huit dispositifs informatiques et de deux unités superfonctionnelles. Toutes les instructions sont exécutées selon le principe SIMD, où une instruction est appliquée à tous les fils d'une chaîne (un terme de l'industrie textile, dans CUDA, il s'agit d'un groupe de 32 threads, la quantité minimale de données traitées par des multiprocesseurs). Cette méthode d'exécution s'appelait SIMT (single instruction multiple threads, une instruction et plusieurs threads).

Chacun des multiprocesseurs dispose de certaines ressources. Il existe donc une mémoire partagée spéciale de 16 kilo-octets par multiprocesseur. Mais ce n'est pas un cache, puisque le programmeur peut l'utiliser pour n'importe quel besoin, comme le Local Store dans le SPU des processeurs Cell. Cette mémoire partagée permet d'échanger des informations entre les threads d'un même bloc. Il est important que tous les threads d'un bloc soient toujours exécutés par le même multiprocesseur. Mais les threads de différents blocs ne peuvent pas échanger de données et vous devez vous rappeler cette limitation. La mémoire partagée est souvent utile, sauf lorsque plusieurs threads accèdent à la même banque de mémoire. Les multiprocesseurs peuvent également accéder à la mémoire vidéo, mais avec des latences plus élevées et un débit inférieur. Pour accélérer l'accès et réduire la fréquence d'accès à la mémoire vidéo, les multiprocesseurs disposent de 8 kilo-octets de cache pour les constantes et les données de texture.

Le multiprocesseur utilise les registres 8192-16384 (pour G8x/G9x et GT2xx, respectivement), communs à tous les threads de tous les blocs exécutés dessus. Le nombre maximum de blocs par multiprocesseur pour G8x/G9x est de huit et le nombre de chaînes est de 24 (768 threads par multiprocesseur). Au total, les cartes vidéo haut de gamme des séries Geforce 8 et 9 peuvent traiter jusqu'à 12 288 threads à la fois. La GeForce GTX 280 basée sur GT200 propose jusqu'à 1024 threads par multiprocesseur, elle dispose de 10 clusters de trois multiprocesseurs traitant jusqu'à 30720 threads. Connaître ces limitations vous permet d'optimiser les algorithmes pour les ressources disponibles.

La première étape lors de la migration d'une application existante vers CUDA consiste à la profiler et à identifier les zones de code qui constituent des goulots d'étranglement qui ralentissent le travail. Si parmi ces sections il y a celles qui conviennent à une exécution parallèle rapide, ces fonctions sont transférées vers les extensions CUDA C pour exécution sur le GPU. Le programme est compilé à l'aide d'un compilateur fourni par Nvidia qui génère du code pour le CPU et le GPU. Lorsqu'un programme est exécuté, le processeur central exécute ses portions de code et le GPU exécute le code CUDA avec les calculs les plus parallèles. Cette partie dédiée au GPU s'appelle le noyau. Le noyau définit les opérations qui seront effectuées sur les données.

La puce vidéo reçoit le noyau et crée des copies pour chaque élément de données. Ces copies sont appelées threads. Un flux contient un compteur, des registres et un état. Pour de gros volumes de données, comme le traitement d’images, des millions de threads sont lancés. Les threads sont exécutés en groupes de 32, appelés warps. Les warps sont exécutés sur des multiprocesseurs de thread spécifiques. Chaque multiprocesseur se compose de huit processeurs de flux cœurs qui exécutent une instruction MAD par cycle d'horloge. Pour exécuter une chaîne à 32 threads, quatre cycles d'horloge du multiprocesseur sont nécessaires (nous parlons de la fréquence du domaine shader, qui est de 1,5 GHz et plus).

Le multiprocesseur n'est pas un processeur multicœur traditionnel ; il est hautement capable de multithread, prenant en charge jusqu'à 32 warps à la fois. À chaque cycle d'horloge, le matériel sélectionne la chaîne à exécuter et passe de l'une à l'autre sans rien perdre. cycles d'horloge. Si nous faisons une analogie avec un processeur central, cela équivaut à l'exécution simultanée de 32 programmes et à la commutation entre eux à chaque cycle d'horloge sans perte de changement de contexte. En réalité, les cœurs de processeur prennent en charge l'exécution simultanée d'un programme et passent aux autres avec un délai de plusieurs centaines de cycles.

Modèle de programmation CUDA

Répétons que CUDA utilise un modèle de calcul parallèle, lorsque chacun des processeurs SIMD exécute en parallèle la même instruction sur différents éléments de données. Le GPU est un appareil informatique, un coprocesseur (appareil) pour le processeur central (hôte), qui possède sa propre mémoire et traite un grand nombre de threads en parallèle. Le noyau est une fonction pour le GPU, exécutée par des threads (une analogie avec les graphiques 3D - un shader).

Nous avons dit plus haut qu'une puce vidéo diffère d'un processeur en ce sens qu'elle peut traiter des dizaines de milliers de threads simultanément, ce qui est typique des graphiques bien parallélisés. Chaque flux est scalaire et ne nécessite pas de regroupement des données dans des vecteurs à 4 composants, ce qui est plus pratique pour la plupart des tâches. Le nombre de threads logiques et de blocs de threads dépasse le nombre de périphériques d'exécution physiques, ce qui offre une bonne évolutivité pour l'ensemble des solutions de l'entreprise.

Le modèle de programmation CUDA implique le regroupement de threads. Les threads sont organisés en blocs de threads : des grilles de threads à une ou deux dimensions qui communiquent entre eux à l'aide de la mémoire partagée et des points de synchronisation. Le programme (noyau) s'exécute sur une grille de blocs de threads, voir la figure ci-dessous. Une grille est exécutée à la fois. Chaque bloc peut avoir une forme à une, deux ou trois dimensions et peut être constitué de 512 threads sur le matériel actuel.

Les blocs de threads sont exécutés en petits groupes appelés warps, dont la taille est de 32 threads. Il s'agit de la quantité minimale de données pouvant être traitées dans des multiprocesseurs. Et comme ce n'est pas toujours pratique, CUDA permet de travailler avec des blocs contenant de 64 à 512 threads.

Regrouper les blocs en grilles vous permet de vous éloigner des restrictions et d'appliquer le noyau à plusieurs threads en un seul appel. Cela aide également à la mise à l’échelle. Si le GPU ne dispose pas de suffisamment de ressources, il exécutera les blocs de manière séquentielle. Sinon, les blocs peuvent être exécutés en parallèle, ce qui est important pour une répartition optimale du travail sur des puces vidéo de différents niveaux, allant du mobile à l'intégré.

Modèle de mémoire CUDA

Le modèle de mémoire dans CUDA se distingue par la possibilité d'adressage d'octets, la prise en charge à la fois de la collecte et de la dispersion. Un assez grand nombre de registres sont disponibles pour chaque processeur de flux, jusqu'à 1024 pièces. L'accès à ceux-ci est très rapide et vous pouvez y stocker des entiers de 32 bits ou des nombres à virgule flottante.

Chaque thread a accès aux types de mémoire suivants :

Mémoire globale la plus grande quantité de mémoire disponible pour tous les multiprocesseurs sur une puce vidéo, la taille varie de 256 Mo à 1,5 Go sur les solutions actuelles (et jusqu'à 4 Go sur Tesla). Il a un débit élevé, supérieur à 100 gigaoctets/s pour les meilleures solutions Nvidia, mais des latences très élevées de plusieurs centaines de cycles. Non mis en cache, prend en charge les instructions génériques de chargement et de stockage, ainsi que les pointeurs de mémoire standard.

Mémoire locale il s'agit d'une petite quantité de mémoire à laquelle un seul processeur de flux a accès. C'est relativement lent - tout comme le système mondial.

La memoire partagée il s'agit d'un bloc mémoire de 16 kilo-octets (en puces vidéo de l'architecture actuelle) avec un accès partagé pour tous les processeurs de flux du multiprocesseur. Cette mémoire est très rapide, tout comme les registres. Il permet aux threads de communiquer, est directement contrôlé par le développeur et a une faible latence. Avantages de la mémoire partagée : utilisation comme cache de premier niveau contrôlé par le programmeur, latence réduite lors de l'accès aux unités d'exécution (ALU) aux données, nombre réduit d'accès à la mémoire globale.

Mémoire constante- une zone mémoire de 64 kilo-octets (la même pour les GPU actuels), en lecture seule par tous les multiprocesseurs. Il est mis en cache à 8 kilo-octets par multiprocesseur. Assez lent - un délai de plusieurs centaines de cycles en l'absence des données nécessaires dans le cache.

Mémoire de textures un bloc de mémoire lisible par tous les multiprocesseurs. L'échantillonnage des données est effectué à l'aide des blocs de texture de la puce vidéo, ce qui permet une interpolation linéaire des données sans frais supplémentaires. 8 kilo-octets sont mis en cache par multiprocesseur. Ralentir au niveau mondial des centaines de cycles de latence lorsqu'il n'y a pas de données dans le cache.

Naturellement, les mémoires globale, locale, de texture et constante sont physiquement la même mémoire, appelée mémoire vidéo locale de la carte vidéo. Leurs différences résident dans les différents algorithmes de mise en cache et modèles d’accès. Le CPU peut uniquement mettre à jour et demander de la mémoire externe : globale, constante et texture.

D'après ce qui a été écrit ci-dessus, il est clair que CUDA suppose une approche particulière du développement, pas tout à fait la même que celle adoptée dans les programmes pour CPU. Vous devez vous rappeler des différents types de mémoire, que la mémoire locale et globale ne sont pas mises en cache et que la latence lors de l'accès est bien supérieure à celle de la mémoire de registre, car elle est physiquement située dans des puces distinctes.

Un modèle de résolution de problèmes typique, mais non obligatoire :

• la tâche est divisée en sous-tâches ;
• les données d'entrée sont divisées en blocs qui s'inscrivent dans la mémoire partagée ;
• chaque bloc est traité par un bloc de threads ;
• le sous-bloc est chargé dans la mémoire partagée à partir de la mémoire globale ;
• des calculs appropriés sont effectués sur les données en mémoire partagée ;
• les résultats sont copiés de la mémoire partagée vers la mémoire globale.

Environnement de programmation

CUDA comprend des bibliothèques d'exécution :

• une partie commune qui fournit des types de vecteurs intégrés et des sous-ensembles d'appels RTL pris en charge sur le CPU et le GPU ;
• Composant CPU pour contrôler un ou plusieurs GPU ;
• Un composant GPU qui fournit des fonctionnalités spécifiques au GPU.

Le processus principal de l'application CUDA s'exécute sur un processeur universel (hôte), il exécute plusieurs copies des processus du noyau sur la carte vidéo. Le code du CPU effectue les opérations suivantes : initialise le GPU, alloue de la mémoire sur la carte vidéo et le système, copie les constantes dans la mémoire de la carte vidéo, lance plusieurs copies des processus du noyau sur la carte vidéo, copie le résultat de la mémoire vidéo, libère le mémoire et sorties.

À titre d'exemple de compréhension, voici le code CPU pour l'addition de vecteurs présenté dans CUDA :

Les fonctions exécutées par la puce vidéo ont les limitations suivantes : il n'y a pas de récursion, pas de variables statiques à l'intérieur des fonctions et pas de nombre variable d'arguments. Deux types de gestion de mémoire sont pris en charge : la mémoire linéaire avec accès via des pointeurs 32 bits et les tableaux CUDA avec accès uniquement via des fonctions de récupération de texture.

Les programmes CUDA peuvent interagir avec les API graphiques : pour restituer les données générées dans le programme, pour lire les résultats du rendu et les traiter à l'aide des outils CUDA (par exemple, lors de la mise en œuvre de filtres de post-traitement). Pour y parvenir, les ressources de l'API graphique peuvent être mappées (en recevant l'adresse de la ressource) dans l'espace mémoire global CUDA. Les types de ressources API graphiques suivants sont pris en charge : objets tampons (PBO/VBO) en OpenGL, tampons de sommets et textures (2D, 3D et cubemaps) Direct3D9.

Étapes de compilation d'une application CUDA :

Les fichiers de code source CUDA C sont compilés à l'aide du programme NVCC, qui est un wrapper sur d'autres outils et les appelle : cudacc, g++, cl, etc. NVCC génère : Du code CPU, qui est compilé avec le reste de l'application écrit en langage pur. C et code objet PTX pour la puce vidéo. Les fichiers exécutables avec du code CUDA nécessitent nécessairement la bibliothèque d'exécution CUDA (cudart) et la bibliothèque principale CUDA (cuda).

Optimisation des programmes sur CUDA

Naturellement, il est impossible d'envisager de sérieux problèmes d'optimisation dans la programmation CUDA dans le cadre d'un article de synthèse. Par conséquent, nous allons simplement parler brièvement des choses de base. Pour utiliser efficacement les capacités de CUDA, vous devez oublier les méthodes habituelles d'écriture de programmes pour le CPU et utiliser des algorithmes bien parallélisés sur des milliers de threads. Il est également important de trouver l'endroit optimal pour stocker les données (registres, mémoire partagée, etc.), de minimiser le transfert de données entre le CPU et le GPU et d'utiliser la mise en mémoire tampon.

En général, lors de l'optimisation d'un programme CUDA, vous devez essayer d'atteindre un équilibre optimal entre la taille et le nombre de blocs. Plus de threads par bloc réduiront l'impact de la latence mémoire, mais réduiront également le nombre de registres disponibles. De plus, un bloc de 512 threads est inefficace ; Nvidia lui-même recommande d'utiliser des blocs de 128 ou 256 threads comme compromis pour obtenir des latences et un nombre de registres optimaux.

Parmi les principaux points d'optimisation des programmes CUDA : l'utilisation la plus active de la mémoire partagée, car elle est beaucoup plus rapide que la mémoire vidéo globale de la carte vidéo ; Les lectures et écritures à partir de la mémoire globale doivent être fusionnées autant que possible. Pour ce faire, vous devez utiliser des types de données spéciaux pour lire et écrire 32/64/128 bits de données à la fois en une seule opération. Si les lectures sont difficiles à combiner, vous pouvez essayer d'utiliser les récupérations de texture.

conclusions

L'architecture logicielle et matérielle présentée par Nvidia pour les calculs sur les puces vidéo CUDA est bien adaptée pour résoudre un large éventail de tâches hautement parallèles. CUDA fonctionne sur une large gamme de GPU Nvidia et améliore le modèle de programmation GPU en le simplifiant considérablement et en ajoutant un grand nombre de fonctionnalités telles que la mémoire partagée, la synchronisation des threads, les calculs en double précision et les opérations sur les nombres entiers.

CUDA est une technologie accessible à tout développeur de logiciels ; elle peut être utilisée par tout programmeur connaissant le langage C. Il suffit de s'habituer aux différents paradigmes de programmation inhérents au calcul parallèle. Mais si l'algorithme est, en principe, bien parallélisé, alors l'étude et le temps consacré à la programmation dans CUDA reviendront plusieurs fois.

Il est probable qu'en raison de l'utilisation généralisée des cartes vidéo dans le monde, le développement du calcul parallèle sur GPU aura un impact considérable sur l'industrie du calcul haute performance. Ces possibilités ont déjà suscité un grand intérêt dans les milieux scientifiques, et pas seulement. Après tout, les opportunités potentielles d'accélérer des dizaines de fois des algorithmes qui peuvent être facilement parallélisés (sur le matériel disponible, ce qui n'est pas moins important) ne sont pas si courantes.

Les processeurs à usage général se développent assez lentement et ne présentent pas de tels sauts de performances. Essentiellement, même si cela semble représenter beaucoup d'argent, toute personne ayant besoin d'une informatique rapide peut désormais disposer d'un superordinateur personnel bon marché sur son bureau, parfois sans même investir davantage, puisque les cartes graphiques Nvidia sont très largement disponibles. Sans parler de l’efficacité accrue en termes de GFLOPS/$ et GFLOPS/W que les fabricants de GPU apprécient tant.

L’avenir de nombreuses technologies informatiques réside clairement dans les algorithmes parallèles, et presque toutes les nouvelles solutions et initiatives vont dans cette direction. Cependant, jusqu'à présent, le développement de nouveaux paradigmes n'en est qu'à ses débuts : vous devez créer manuellement des threads et planifier l'accès à la mémoire, ce qui complique les tâches par rapport à la programmation conventionnelle. Mais la technologie CUDA a fait un pas dans la bonne direction et elle apparaît clairement comme une solution réussie, surtout si Nvidia parvient à convaincre un maximum de développeurs de ses avantages et de ses perspectives.

Mais bien entendu, les GPU ne remplaceront pas les CPU. Dans leur forme actuelle, ils ne sont pas destinés à cela. Alors que les puces vidéo s'orientent progressivement vers le CPU, devenant de plus en plus universelles (calculs à virgule flottante simple et double précision, calculs d'entiers), les CPU deviennent de plus en plus « parallèles », se dotant d'un grand nombre de cœurs, de technologies multi-threading. , sans compter l'apparition de blocs SIMD et de projets de processeurs hétérogènes. Très probablement, le GPU et le CPU fusionneront simplement à l'avenir. On sait que de nombreuses entreprises, dont Intel et AMD, travaillent sur des projets similaires. Et peu importe que les GPU soient absorbés par le CPU, ou vice versa.

Dans l'article, nous avons principalement parlé des avantages de CUDA. Mais il y a aussi un problème. L’un des rares inconvénients de CUDA est sa mauvaise portabilité. Cette architecture ne fonctionne que sur les puces vidéo de cette société, et pas sur toutes, mais à commencer par les séries Geforce 8 et 9 et les Quadro et Tesla correspondantes. Oui, il existe de nombreuses solutions de ce type dans le monde ; Nvidia cite un chiffre de 90 millions de puces vidéo compatibles CUDA. C'est tout simplement génial, mais les concurrents proposent leurs propres solutions différentes de CUDA. Ainsi, AMD a Stream Computing, Intel aura Ct à l'avenir.

Quelle technologie gagnera, se généralisera et vivra plus longtemps que les autres - seul le temps nous le dira. Mais CUDA a de bonnes chances, car comparé à Stream Computing, par exemple, il représente un environnement de programmation plus développé et plus facile à utiliser dans le langage C classique. Peut-être qu'un tiers peut aider à la détermination en proposant une solution générale. Par exemple, dans la prochaine mise à jour de DirectX sous la version 11, Microsoft a promis des shaders informatiques, qui pourraient devenir une sorte de solution moyenne qui conviendrait à tout le monde, ou presque.

À en juger par les données préliminaires, ce nouveau type de shader emprunte beaucoup au modèle CUDA. Et en programmant dès maintenant dans cet environnement, vous pouvez obtenir des avantages immédiats et les compétences nécessaires pour l’avenir. Du point de vue du calcul haute performance, DirectX présente également le net inconvénient d'une mauvaise portabilité, car l'API est limitée à la plate-forme Windows. Cependant, un autre standard est en cours de développement : l'initiative multiplateforme ouverte OpenCL, soutenue par la plupart des entreprises, notamment Nvidia, AMD, Intel, IBM et bien d'autres.

N'oubliez pas que le prochain article sur CUDA explorera des applications pratiques spécifiques du calcul scientifique et autre calcul non graphique réalisées par des développeurs de différentes parties de notre planète utilisant Nvidia CUDA.

Noyaux CUDA - symbole unités de calcul scalaires en puces vidéo NVIDIA, commençant par G80 (GeForce 8 xxx, Tesla CD-S870, FX4/5600 , 360M). Les puces elles-mêmes sont des dérivés de l'architecture. D'ailleurs, parce que l'entreprise NVIDIA si volontiers entrepris le développement de ses propres processeurs Série Tégra, également basé sur RISQUE architecture. J'ai beaucoup d'expérience avec ces architectures.

CUDA le noyau en contient un un vecteur Et un scalaire unités qui effectuent une opération vectorielle et une opération scalaire par cycle d'horloge, transférant les calculs vers un autre multiprocesseur ou vers un autre pour un traitement ultérieur. Un ensemble de centaines et de milliers de ces cœurs représente une puissance de calcul importante et peut effectuer diverses tâches en fonction des besoins, à condition qu'il existe certains logiciels de support. Application peuvent être variés : décodage de flux vidéo, accélération graphique 2D/3D, cloud computing, analyses mathématiques spécialisées, etc.

Très souvent, combinés Cartes professionnelles NVidia Tesla Et NVidia Quadro, constituent l’épine dorsale des supercalculateurs modernes.

CUDA— les amandes n'ont pas subi de modifications significatives depuis G80, mais leur nombre augmente (avec d'autres blocs - ROP, Unités de texture& etc) et l'efficacité des interactions parallèles entre elles (les modules sont améliorés Giga-fil).

Par exemple:

GeForce

GTX 460 – 336 cœurs CUDA

GTX 580 – 512 cœurs CUDA

8800GTX – 128 cœurs CUDA

Du nombre de processeurs de flux ( CUDA), les performances dans les calculs de shaders augmentent presque proportionnellement (avec une augmentation uniforme du nombre d'autres éléments).

À partir de la puce GK110(NVidiaGeForce GTX 680) - CUDA les cœurs n'ont plus une double fréquence, mais une fréquence commune à tous les autres blocs de puces. Au lieu de cela, leur nombre a augmenté d'environ trois fois par rapport à la génération précédente G110.

La nouvelle technologie est comme une espèce évolutive nouvellement émergée. Une créature étrange, contrairement à de nombreux anciens. Parfois gênant, parfois drôle. Et de prime abord, ses nouvelles qualités ne semblent aucunement adaptées à ce monde posé et stable.

Cependant, un peu de temps passe et il s'avère que le débutant court plus vite, saute plus haut et est 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 sert à rien de se disputer avec cet ancien maladroit. Il vaut mieux être ami avec lui, et encore mieux organiser une symbiose. Vous verrez qu'il y aura plus de mouches.

La technologie GPGPU (General-Purpose Graphics Processing Units - processeur graphique à usage général) n'a longtemps existé que dans les calculs théoriques d'universitaires intelligents. Sinon comment? Proposer de changer radicalement le processus informatique développé au fil des décennies en confiant le calcul de ses branches parallèles à une carte vidéo, seuls les théoriciens en sont capables.

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

Mais la technologie GPGPU n’allait pas rester longtemps dans les pages des revues universitaires. Après avoir gonflé 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 processeurs graphiques GeForce fabriqués par nVidia.

Grâce à CUDA, les technologies GPGPU sont devenues courantes. Et maintenant, seuls les développeurs de systèmes de programmation les plus myopes et couverts d'une épaisse couche de paresse ne déclarent pas de support pour CUDA avec leur produit. Les publications informatiques ont considéré comme un honneur de présenter les détails de la technologie dans de nombreux articles scientifiques de vulgarisation, et les concurrents se sont immédiatement mis à la recherche de modèles et de compilateurs croisés pour développer quelque chose de similaire.

La reconnaissance publique est un rêve non seulement pour les starlettes en herbe, mais aussi pour les technologies nouvellement nées. Et CUDA a eu de la chance. Elle est bien connue, on parle et on écrit sur elle.

Ils écrivent simplement comme s’ils continuaient à discuter du GPGPU dans de grosses revues scientifiques. Ils bombardent le lecteur avec un tas de termes comme « grille », « SIMD », « distorsion », « hôte », « texture et mémoire constante ». Ils le plongent tout en haut dans les organigrammes des GPU nVidia, le conduisent sur des chemins sinueux d'algorithmes parallèles et (le geste le plus fort) affichent de longues listes de codes en langage C. En conséquence, il s'avère qu'à l'entrée de l'article nous avons un nouveau lecteur avec un désir ardent de comprendre CUDA, et à la sortie nous avons le même lecteur, mais avec une tête gonflée remplie d'un fouillis de faits, de diagrammes , code, algorithmes et termes.

Pendant ce temps, le but de toute technologie est de nous faciliter la vie. Et CUDA fait un excellent travail avec ça. Les résultats de son travail convaincront tout sceptique mieux que des centaines de schémas et d’algorithmes.

Pas partout

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

Et pourtant, avant d’examiner les résultats des travaux de CUDA visant à faciliter la vie de l’utilisateur moyen, il convient d’en comprendre toutes les limites. Comme avec un génie : n'importe quel désir, mais un. CUDA a aussi son talon d'Achille. L’un d’eux concerne les limitations des plates-formes sur lesquelles il peut fonctionner.

La liste des cartes vidéo nVidia prenant en charge CUDA est présentée dans une liste spéciale appelée Produits compatibles CUDA. La liste est assez impressionnante, mais facile à classer. Le support de CUDA n'est pas refusé :

Modèles nVidia GeForce des séries 8e, 9e, 100e, 200e et 400e avec un minimum de 256 Mo de mémoire vidéo intégrée. La prise en charge s'étend aux cartes de bureau et mobiles.

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

Toutes les solutions de la série de netbooks nvidia ION.

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

Par conséquent, avant d'utiliser des produits logiciels basés sur CUDA, il convient de consulter cette liste de favoris.

En plus de la carte vidéo elle-même, un pilote approprié est requis pour prendre en charge CUDA. C'est le lien entre les processeurs centraux et graphiques, agissant comme une sorte d'interface logicielle pour accéder au code du programme et aux données du trésor multicœur du GPU. Pour vous assurer de ne pas commettre d'erreur, nVidia vous recommande de visiter la page des pilotes et d'obtenir la dernière version.

... mais le processus lui-même

Comment fonctionne CUDA ? Comment expliquer le processus complexe de calcul parallèle sur une architecture matérielle GPU particulière sans plonger le lecteur dans l'abîme des termes spécifiques ?

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

Sur le plan architectural, l'unité centrale de traitement (CPU) et son homologue graphique (GPU) sont conçus différemment. Si l'on fait une analogie avec le monde de l'industrie automobile, alors le CPU est un break, un de ceux qu'on appelle une « grange ». Cela ressemble à une voiture de tourisme, mais en même temps (du point de vue des développeurs) « c’est un Suisse, un faucheur et un joueur de pipe ». Remplit à la fois le rôle d’un petit camion, d’un bus et d’une berline hypertrophiée. Break, en somme. Il comporte peu de cœurs de cylindre, mais ils gèrent presque toutes les tâches et l'impressionnante mémoire cache est capable de stocker un tas de données.

Mais le GPU est une voiture de sport. Il n’y a qu’une seule fonction : amener le pilote à la ligne d’arrivée le plus rapidement possible. Par conséquent, pas de grande mémoire de coffre, pas de sièges supplémentaires. Mais il y a des centaines de fois plus de cœurs de cylindre que de processeur.

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

Contrairement au processeur central, qui est 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'une tâche : transformer un tas de polygones à l'entrée en un tas de pixels à le résultat. De plus, ce problème peut être résolu en parallèle en utilisant des centaines de cœurs de calcul relativement simples dans le GPU.

Alors, quel genre de tandem peut-il y avoir entre un break et une voiture de sport ? CUDA fonctionne à peu près comme ceci : le programme s'exécute sur le CPU jusqu'à ce qu'il y ait une section de code qui peut être exécutée en parallèle. Ensuite, au lieu d’être exécuté 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.

Technologiquement, rien ne change pour le programmeur. Le code des programmes CUDA est écrit en langage C. Plus précisément, dans son dialecte spécial « C with streams » (C with streams). Développée à Stanford, cette extension du langage C s'appelle Brook. L'interface qui transfère le 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 afin que, pour le programmeur, le GPU ressemble à un coprocesseur CPU. Très similaire à l’utilisation d’un coprocesseur mathématique aux débuts de l’informatique personnelle. Avec l'avènement de Brook, des cartes vidéo prenant en charge CUDA et de leurs pilotes, tout programmeur est désormais en mesure d'accéder au GPU dans ses programmes. Mais avant, ce chamanisme appartenait à un cercle restreint de personnes sélectionnées qui passaient des années à perfectionner les techniques de programmation pour les moteurs graphiques DirectX ou OpenGL.

Dans le tonneau de ce miel prétentieux - les louanges de CUDA - il vaut la peine de mettre un frein à la pommade, c'est-à-dire des restrictions. Tous les problèmes qui doivent être programmés ne peuvent pas être résolus à l'aide de CUDA. Il ne sera pas possible d'accélérer la résolution des tâches de bureau courantes, mais vous pouvez faire confiance à CUDA pour calculer le comportement de milliers de combattants du même type dans World of Warcraft. Mais c’est une tâche inventée. Regardons des exemples de ce que CUDA résout déjà de manière très efficace.

Œuvres justes

CUDA est une technologie très pragmatique. Après avoir implémenté son support dans ses cartes vidéo, nVidia s'attendait à juste titre à ce que la bannière CUDA soit reprise par de nombreux passionnés tant dans le milieu universitaire que dans le commerce. Et c’est ce qui s’est passé. Les projets basés sur CUDA vivent et apportent des avantages.

NVIDIA PhysX

Lorsqu'ils font la publicité de leur prochain chef-d'œuvre de jeu, les fabricants mettent souvent l'accent sur son réalisme 3D. Mais aussi réel que soit le monde du jeu 3D, si les lois élémentaires de la physique, telles que la gravité, la friction et l'hydrodynamique, ne sont pas mises en œuvre correctement, le mensonge se fera instantanément sentir.

L'une des capacités du moteur physique NVIDIA PhysX est le travail réaliste avec les tissus.

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

Cela est devenu possible grâce aux statistiques. Il a été statistiquement prouvé que, quelle que soit la complexité du rendu d'un GPU, certains de ses cœurs restent inactifs. C'est sur ces cœurs que tourne le moteur PhysX.

Grâce à CUDA, la part du lion des calculs liés à la physique du monde du jeu a commencé à être effectuée sur la 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 dans le gameplay avec PhysX fonctionnant sur CUDA a augmenté d'au moins un ordre de grandeur. La probabilité de réaliser les lois physiques a également augmenté. CUDA s'occupe du calcul de routine de la mise en œuvre du frottement, de la gravité et d'autres éléments qui nous sont familiers pour les objets multidimensionnels. Désormais, non seulement les héros et leur équipement s'intègrent parfaitement aux lois du monde physique que nous connaissons, mais aussi la poussière, le brouillard, l'onde de souffle, les flammes et l'eau.

Version CUDA du package de compression de texture NVIDIA Texture Tools 2

Aimez-vous les objets réalistes dans les jeux modernes ? Cela vaut la peine de remercier les développeurs de textures. Mais plus la texture est réelle, plus son volume est grand. Plus cela occupe une mémoire précieuse. Pour éviter cela, les textures sont pré-compressées et décompressées dynamiquement selon les besoins. Et la compression et la décompression sont de purs calculs. Pour travailler avec des textures, nVidia a publié le package NVIDIA Texture Tools. Il prend en charge la compression et la décompression efficaces des textures DirectX (le format dit HF). 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 inclut la prise en charge de CUDA. Selon nVidia, cela permet de multiplier par 12 les performances dans les tâches de compression et de décompression de texture. Cela signifie que les images de jeu 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. Le gain de performances lors de la compression et de la décompression des textures est évident.

L'utilisation de CUDA peut améliorer considérablement l'efficacité de la vidéosurveillance.

Traitement du flux vidéo en temps réel

Quoi qu’on en dise, le monde actuel, du point de vue de l’espionnage, est bien plus proche du monde du Big Brother d’Orwell qu’il n’y paraît. Les automobilistes et les visiteurs des lieux publics ressentent le regard des caméras vidéo.

Des fleuves d'informations vidéo se déversent dans les centres de leur traitement et... se jettent dans un lien étroit - une personne. Dans la plupart des cas, il est la dernière autorité à surveiller le monde de la vidéo. De plus, l’autorité n’est pas la plus efficace. Cligne des yeux, se laisse distraire 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 atteint une vitesse de 30 ips. Par rapport à la mise en œuvre d'un tel algorithme sur les processeurs multicœurs modernes, CUDA offre des performances deux ou trois fois supérieures, et cela, voyez-vous, c'est beaucoup.

Conversion vidéo, filtrage audio

Le convertisseur vidéo Badaboom est le premier à utiliser CUDA pour accélérer la conversion.

C'est agréable de regarder un nouveau produit de location vidéo 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 est compréhensible, 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. Le premier signe est le convertisseur Badaboom d'Elevental. Les développeurs de Badaboom ont pris la bonne décision en choisissant CUDA. Les tests montrent qu'il convertit un film standard d'une heure et demie au format iPhone/iPod Touch en moins de vingt minutes. Et ceci 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 donnerait un demi-royaume pour un crossover FIR efficace - un ensemble de filtres qui divisent le spectre sonore en plusieurs bandes. Ce processus demande beaucoup de travail et, avec un volume important de matériel audio, oblige l'ingénieur du son à « fumer » pendant plusieurs heures. La mise en œuvre d'un crossover FIR basé sur CUDA accélère son fonctionnement des centaines de fois.

L'avenir de CUDA

Ayant fait de la technologie GPGPU une réalité, CUDA ne se repose pas sur ses lauriers. Comme cela arrive partout, le principe de réflexion fonctionne dans CUDA : désormais non seulement l'architecture des processeurs vidéo nVidia influence le développement des versions du SDK CUDA, mais la technologie CUDA elle-même oblige nVidia à reconsidérer l'architecture de ses puces. Un exemple d’une telle réflexion est la plateforme nVidia ION. Sa deuxième version est spécialement optimisée pour résoudre les problèmes CUDA. Cela signifie que même avec des solutions matérielles relativement peu coûteuses, les consommateurs bénéficieront de toute la puissance et des capacités brillantes de CUDA.

Technologie CUDA

Vladimir Frolov,[email protégé]

annotation

L'article parle de la technologie CUDA, qui permet à un programmeur d'utiliser des cartes vidéo comme unités informatiques puissantes. Les outils fournis par Nvidia permettent d'écrire des programmes d'unité de traitement graphique (GPU) dans un sous-ensemble du langage C++. Cela évite au programmeur d'avoir à utiliser des shaders et à comprendre le fonctionnement du pipeline graphique. L'article fournit des exemples de programmation utilisant CUDA et diverses techniques d'optimisation.

1. Introduction

Le développement de la technologie informatique a progressé rapidement au cours des dernières décennies. Si vite que les développeurs de processeurs ont déjà presque atteint ce qu’on appelle « l’impasse du silicium ». Une augmentation effrénée de la fréquence d'horloge est devenue impossible pour un certain nombre de raisons technologiques graves.

C'est en partie pourquoi tous les fabricants de systèmes informatiques modernes s'orientent vers une augmentation du nombre de processeurs et de cœurs, plutôt que d'augmenter la fréquence d'un processeur. Le nombre de cœurs d’unité centrale (CPU) dans les systèmes avancés est désormais déjà de 8.

Une autre raison est la vitesse relativement faible de la RAM. Quelle que soit la vitesse de fonctionnement du processeur, les goulots d'étranglement, comme le montre la pratique, ne sont pas du tout des opérations arithmétiques, mais plutôt des accès infructueux à la mémoire - des échecs de cache.

Cependant, si l’on se tourne vers les processeurs graphiques GPU (Graphics Processing Unit), ils ont emprunté la voie du parallélisme bien plus tôt. Dans les cartes vidéo actuelles, par exemple dans la GF8800GTX, le nombre de processeurs peut atteindre 128. Les performances de tels systèmes, avec une programmation habile, peuvent être assez importantes (Fig. 1).

Riz. 1. Nombre d'opérations en virgule flottante pour CPU et GPU

Lorsque les premières cartes vidéo sont apparues pour la première fois sur le marché, il s'agissait de dispositifs assez simples (par rapport au processeur central) hautement spécialisés, conçus pour soulager le processeur du fardeau de la visualisation des données bidimensionnelles. Avec le développement de l'industrie du jeu et l'émergence de jeux en trois dimensions tels que Doom (Fig. 2) et Wolfenstein 3D (Fig. 3), le besoin de visualisation 3D s'est fait sentir.

Chiffres 2,3. Jeux Doom et Wolfenstein 3D

Depuis le moment où 3Dfx a créé les premières cartes vidéo Voodoo (1996) jusqu'en 2001, seul un ensemble fixe d'opérations sur les données d'entrée a été implémenté dans le GPU.

Les programmeurs n'avaient pas le choix dans l'algorithme de rendu, et pour augmenter la flexibilité, des shaders sont apparus - de petits programmes exécutés par la carte vidéo pour chaque sommet ou pour chaque pixel. Leurs tâches comprenaient des transformations sur les sommets et des ombrages, en calculant par exemple l'éclairage en un point, à l'aide du modèle Phong.

Bien que les shaders aient parcouru un long chemin ces jours-ci, il faut comprendre qu'ils ont été conçus pour des tâches hautement spécialisées telles que la transformation 3D et la rastérisation. Alors que les GPU évoluent vers des systèmes multiprocesseurs généralistes, les langages shader restent très spécialisés.

Ils peuvent être comparés à FORTRAN dans le sens où, comme FORTRAN, ils ont été les premiers, mais conçus pour résoudre un seul type de problème. Les shaders sont peu utiles pour résoudre d'autres problèmes que les transformations 3D et la rastérisation, tout comme FORTRAN n'est pas pratique pour résoudre des problèmes non liés aux calculs numériques.

Aujourd'hui, on a tendance à utiliser les cartes vidéo de manière non conventionnelle pour résoudre des problèmes dans les domaines de la mécanique quantique, de l'intelligence artificielle, des calculs physiques, de la cryptographie, de la visualisation physiquement correcte, de la reconstruction à partir de photographies, de la reconnaissance, etc. Ces tâches sont peu pratiques à résoudre dans le cadre des API graphiques (DirectX, OpenGL), puisque ces API ont été créées pour des applications complètement différentes.

Le développement de la programmation générale sur GPU (General Programming on GPU, GPGPU) a logiquement conduit à l'émergence de technologies visant un éventail de tâches plus large que la rastérisation. En conséquence, Nvidia a créé Compute Unified Device Architecture (ou CUDA en abrégé) et la société concurrente ATI a créé la technologie STREAM.

Il convient de noter qu'au moment de la rédaction de cet article, la technologie STREAM était loin derrière CUDA en termes de développement, et elle ne sera donc pas prise en compte ici. Nous nous concentrerons sur CUDA, une technologie GPGPU qui permet d'écrire des programmes dans un sous-ensemble du langage C++.

2. Différence fondamentale entre CPU et GPU

Jetons un coup d'œil rapide à certaines des différences significatives entre les domaines et les fonctionnalités des applications CPU et carte vidéo.

2.1. Possibilités

Le processeur est initialement conçu pour résoudre des problèmes généraux et fonctionne avec une mémoire adressable de manière aléatoire. Les programmes sur le CPU peuvent accéder directement à toutes les cellules mémoire linéaires et homogènes.

Ce n'est pas le cas des GPU. Comme vous l'apprendrez en lisant cet article, CUDA possède jusqu'à 6 types de mémoire. Vous pouvez lire à partir de n’importe quelle cellule physiquement accessible, mais vous ne pouvez pas écrire dans toutes les cellules. La raison en est que le GPU est de toute façon un appareil spécifique conçu à des fins spécifiques. Cette limitation a été introduite pour augmenter la vitesse de certains algorithmes et réduire le coût des équipements.

2.2. Performances de la mémoire

Un problème récurrent avec la plupart des systèmes informatiques est que la mémoire est plus lente que le processeur. Les fabricants de processeurs résolvent ce problème en introduisant des caches. Les zones de mémoire les plus fréquemment utilisées sont placées en mémoire ou en mémoire cache, fonctionnant à la fréquence du processeur. Cela vous permet de gagner du temps lors de l'accès aux données les plus fréquemment utilisées et de charger le processeur avec les calculs réels.

Notez que les caches sont essentiellement transparents pour le programmeur. Tant en lecture qu'en écriture, les données ne vont pas directement dans la RAM, mais passent par des caches. Cela permet notamment de lire rapidement une valeur immédiatement après son écriture.

Les GPU (nous entendons ici les cartes vidéo GF huitième série) ont également des caches, et ils sont également importants, mais ce mécanisme n'est pas aussi puissant que sur un CPU. Premièrement, tous les types de mémoire ne sont pas mis en cache et, deuxièmement, les caches sont en lecture seule.

Sur les GPU, les accès mémoire lents sont masqués grâce au calcul parallèle. Pendant que certaines tâches attendent des données, d’autres travaillent, prêtes pour les calculs. C’est l’un des principes fondamentaux de CUDA, qui peut grandement améliorer les performances du système dans son ensemble.

3. Noyau CUDA

3.1. Modèle de diffusion en continu

L'architecture informatique CUDA est basée sur le conceptune commande pour plusieurs données(Single Instruction Multiple Data, SIMD) et concept multiprocesseur.

Le concept SIMD signifie qu'une instruction peut traiter plusieurs données simultanément. Par exemple, la commande addps sur les processeurs Pentium 3 et plus récents vous permet d'ajouter 4 nombres à virgule flottante simple précision à la fois.

Un multiprocesseur est un processeur SIMD multicœur qui permet d'exécuter une seule instruction sur tous les cœurs à un moment donné. Chaque cœur multiprocesseur est scalaire, c'est-à-dire il ne prend pas en charge les opérations vectorielles pures.

Avant de continuer, introduisons quelques définitions. Veuillez noter que dans cet article, nous désignerons le périphérique et l'hôte d'une manière complètement différente de celle à laquelle la plupart des programmeurs sont habitués. Nous utiliserons ces termes pour éviter les divergences avec la documentation CUDA.

Dans notre article, par appareil, nous entendrons un adaptateur vidéo prenant en charge le pilote CUDA, ou un autre appareil spécialisé conçu pour exécuter des programmes utilisant CUDA (tels que NVIDIA Tesla). Dans notre article, nous considérerons le GPU uniquement comme un périphérique logique, en évitant les détails spécifiques d'implémentation.

Nous appellerons un hôte un programme dans la RAM normale de l'ordinateur qui utilise le processeur et exécute des fonctions de contrôle pour travailler avec l'appareil.

En fait, la partie de votre programme qui s'exécute sur le CPU est hôte, et votre carte vidéo - appareil. Logiquement, l'appareil peut être représenté comme un ensemble de multiprocesseurs (Fig. 4) plus un pilote CUDA.

Riz. 4. Appareil

Supposons que nous souhaitions exécuter une certaine procédure sur notre appareil dans N threads (c'est-à-dire que nous souhaitons paralléliser son travail). D'après la documentation CUDA, appelons cette procédure le noyau.

Une caractéristique de l'architecture CUDA est son organisation en grille de blocs, ce qui est inhabituel pour les applications multithread (Fig. 5). Dans ce cas, le pilote CUDA distribue indépendamment les ressources du périphérique entre les threads.

Riz. 5. Organisation des fils de discussion

En figue. 5. Le noyau est désigné comme Kernel. Tous les threads exécutant ce noyau sont combinés en blocs (Block), et les blocs, à leur tour, sont combinés dans une grille (Grid).

Comme le montre la figure 5, des indices bidimensionnels sont utilisés pour identifier les threads. Les développeurs CUDA ont offert la possibilité de travailler avec des index tridimensionnels, bidimensionnels ou simples (unidimensionnels), selon ce qui convient le mieux au programmeur.

En général, les indices sont des vecteurs tridimensionnels. Pour chaque thread, on connaîtra : l'index du thread à l'intérieur du bloc threadIdx et l'index du bloc à l'intérieur de la grille blockIdx. Au démarrage, tous les threads ne différeront que par ces index. En fait, c'est grâce à ces index que le programmeur exerce un contrôle, déterminant quelle partie de ses données est traitée dans chaque thread.

La réponse à la question de savoir pourquoi les développeurs ont choisi cette organisation particulière n'est pas triviale. L'une des raisons est qu'un bloc est assuré d'être exécuté sur une périphérique multiprocesseur, mais un multiprocesseur peut exécuter plusieurs blocs différents. Les autres raisons apparaîtront clairement plus loin dans l’article.

Un bloc de tâches (threads) est exécuté sur un multiprocesseur en parties, ou pools, appelés warps. La taille de chaîne actuelle des cartes vidéo prenant en charge CUDA est de 32 threads. Les tâches à l'intérieur du pool de distorsion sont exécutées dans le style SIMD, c'est-à-dire Tous les threads de Warp ne peuvent exécuter qu'une seule instruction à la fois.

Une mise en garde doit être faite ici. Dans les architectures modernes au moment de la rédaction de cet article, le nombre de processeurs à l'intérieur d'un multiprocesseur est de 8 et non de 32. Il s'ensuit que la totalité du warp n'est pas exécutée simultanément, elle est divisée en 4 parties, qui sont exécutées séquentiellement (puisque les processeurs sont scalaires).

Mais premièrement, les développeurs CUDA ne réglementent pas strictement la taille de la chaîne. Dans leurs travaux, ils mentionnent le paramètre de taille de chaîne, et non le nombre 32. Deuxièmement, d'un point de vue logique, la chaîne est l'union minimale de threads dont on peut dire que tous les threads qu'il contient sont exécutés simultanément - et au même moment. en même temps, il n'y a aucune hypothèse concernant le reste, le système ne sera pas réalisé.

3.1.1. Ramification

La question se pose immédiatement : si au même moment tous les threads à l’intérieur d’un warp exécutent la même instruction, alors qu’en est-il des branches ? Après tout, s'il y a une branche dans le code du programme, les instructions seront différentes. Une solution standard pour la programmation SIMD est utilisée ici (Figure 6).

Riz. 6. Organisation du branchement en SIMD

Disons que nous avons le code suivant :

si (suite)B ;

Dans le cas de SISD (Single Instruction Single Data), nous exécutons l'instruction A, vérifions la condition, puis exécutons les instructions B et D (si la condition est vraie).

Disons maintenant 10 threads exécutés dans le style SIMD. Dans les 10 threads, nous exécutons l'instruction A, puis nous vérifions la condition cond et il s'avère que dans 9 threads sur 10, c'est vrai, et dans un thread, c'est faux.

Il est clair que nous ne pouvons pas lancer 9 threads pour exécuter l'opérateur B, et le reste pour exécuter l'opérateur C, car une seule instruction peut être exécutée dans tous les threads en même temps. Dans ce cas, vous devez faire ceci : d'abord, nous « tuons » le thread séparable afin qu'il ne gâche les données de personne, et exécutons les 9 threads restants. Ensuite, nous « tuons » 9 threads qui ont exécuté l'opérateur B et passons par un thread avec l'opérateur C. Après cela, les threads fusionnent à nouveau et exécutent l'opérateur D en même temps.

Le triste résultat est que non seulement les ressources du processeur sont gaspillées en broyant des bits vides dans les threads divisés, mais ce qui est bien pire est que nous finirons par devoir exécuter les DEUX branches.

Cependant, tout n’est pas aussi mauvais qu’il y paraît à première vue. Un très gros avantage de la technologie est que ces astuces sont exécutées dynamiquement par le pilote CUDA et sont totalement transparentes pour le programmeur. Dans le même temps, lorsqu'il s'agit de commandes SSE de processeurs modernes (en particulier dans le cas d'une tentative d'exécution simultanée de 4 copies de l'algorithme), le programmeur lui-même doit s'occuper des détails : combiner les données en quads, ne pas oublier l'alignement. , et généralement écrire à bas niveau, en fait, comment en assembleur.

De tout ce qui précède, une conclusion très importante découle. Les branches n’entraînent pas de dégradation des performances en elles-mêmes. Les seules branches nuisibles sont celles où les threads divergent au sein du même pool de threads de distorsion. De plus, si les threads divergent au sein d'un même bloc, mais dans des pools de distorsion différents, ou au sein de différents blocs, cela n'a absolument aucun effet.

3.1.2. Communication entre les threads

Au moment de la rédaction de cet article, toute interaction entre threads (synchronisation et échange de données) n’était possible qu’au sein d’un bloc. Autrement dit, il est impossible d'organiser l'interaction entre les threads de différents blocs en utilisant uniquement des capacités documentées.

Quant aux fonctionnalités non documentées, leur utilisation est fortement déconseillée. La raison en est qu’ils s’appuient sur les fonctionnalités matérielles spécifiques d’un système particulier.

La synchronisation de toutes les tâches au sein d'un bloc s'effectue en appelant la fonction __synchtreads. L'échange de données est possible via la mémoire partagée, puisqu'elle est commune à toutes les tâches du bloc.

3.2. Mémoire

CUDA distingue six types de mémoire (Fig. 7). Il s'agit de registres, de mémoire locale, globale, partagée, constante et de texture.

Cette abondance est due aux spécificités de la carte vidéo et à son objectif principal, ainsi qu'à la volonté des développeurs de rendre le système aussi bon marché que possible, sacrifiant dans divers cas soit la polyvalence, soit la vitesse.

Riz. 7. Types de mémoire dans CUDA

3.2.0. Registres

Dans la mesure du possible, le compilateur essaie de placer toutes les variables de fonction locales dans des registres. Ces variables sont accessibles à vitesse maximale. Dans l'architecture actuelle, 8 192 registres 32 bits sont disponibles par multiprocesseur. Afin de déterminer combien de registres sont disponibles pour un thread, vous devez diviser ce nombre (8192) par la taille du bloc (le nombre de threads qu'il contient).

Avec la division habituelle de 64 threads par bloc, le résultat n'est que de 128 registres (il existe quelques critères objectifs, mais 64 conviennent en moyenne pour de nombreuses tâches). En réalité, nvcc n'attribuera jamais 128 registres. Habituellement, cela ne donne pas plus de 40 et les variables restantes iront dans la mémoire locale. Cela se produit parce que plusieurs blocs peuvent être exécutés sur un même multiprocesseur. Le compilateur essaie de maximiser le nombre de blocs travaillant simultanément. Pour plus d'efficacité, vous devriez essayer d'occuper moins de 32 registres. Ensuite, théoriquement, 4 blocs (8 warps, si 64 threads dans un bloc) peuvent être lancés sur un multiprocesseur. Cependant, ici, vous devez également prendre en compte la quantité de mémoire partagée occupée par les threads, car si un bloc occupe toute la mémoire partagée, deux de ces blocs ne peuvent pas être exécutés simultanément sur un multiprocesseur.

3.2.1. Mémoire locale

Dans les cas où les données de la procédure locale sont trop volumineuses ou si le compilateur ne peut pas calculer une étape d'accès constante pour celles-ci, il peut les placer dans la mémoire locale. Cela peut être facilité, par exemple, en créant des pointeurs pour des types de tailles différentes.

Physiquement, la mémoire locale est analogue à la mémoire globale et fonctionne à la même vitesse. Au moment de la rédaction de cet article, il n'existe aucun mécanisme empêchant explicitement le compilateur d'utiliser la mémoire locale pour des variables spécifiques. Comme il est assez difficile de contrôler la mémoire locale, il vaut mieux ne pas l'utiliser du tout (voir section 4 « Recommandations d'optimisation »).

3.2.2. Mémoire globale

Dans la documentation CUDA comme l'une des principales réalisationsLa technologie offre la possibilité d'un adressage arbitraire de la mémoire globale. Autrement dit, vous pouvez lire à partir de n’importe quelle cellule mémoire et vous pouvez également écrire dans une cellule arbitraire (ce n’est généralement pas le cas sur un GPU).

Cependant, dans ce cas, la polyvalence doit être payée par la rapidité. La mémoire globale n'est pas mise en cache. Cela fonctionne très lentement, le nombre d'accès à la mémoire globale doit dans tous les cas être minimisé.

La mémoire globale est principalement nécessaire pour sauvegarder les résultats d'un programme avant de les envoyer à l'hôte (en DRAM standard). La raison en est que la mémoire globale est le seul type de mémoire dans laquelle vous pouvez écrire n'importe quoi.

Les variables déclarées avec le qualificatif __global__ sont allouées à la mémoire globale. La mémoire globale peut également être allouée dynamiquement en appelant cudaMalloc(void* mem, int size) sur l'hôte. Cette fonction ne peut pas être appelée depuis l'appareil. Il s'ensuit que l'allocation de mémoire doit être gérée par le programme hôte exécuté sur le processeur. Les données de l'hôte peuvent être envoyées à l'appareil en appelant la fonction cudaMemcpy :

cudaMemcpy (void* gpu_mem, void* cpu_mem, taille int, cudaMemcpyHostToDevice);

Vous pouvez faire la procédure inverse exactement de la même manière :

cudaMemcpy(void* cpu_mem, void* gpu_mem, taille int, cudaMemcpyDeviceToHost);

Cet appel est également effectué depuis l'hôte.

Lorsque vous travaillez avec la mémoire globale, il est important de suivre la règle de fusion. L'idée principale est que les threads doivent accéder à des cellules mémoire consécutives, 4, 8 ou 16 octets. Dans ce cas, le tout premier thread doit accéder à une adresse alignée respectivement sur une limite de 4, 8 ou 16 octets. Les adresses renvoyées par cudaMalloc sont alignées sur une limite d'au moins 256 octets.

3.2.3. La memoire partagée

La mémoire partagée est une mémoire non mise en cache mais rapide. Il est recommandé de l'utiliser comme cache géré. Seuls 16 Ko de mémoire partagée sont disponibles par multiprocesseur. En divisant ce nombre par le nombre de tâches dans le bloc, nous obtenons la quantité maximale de mémoire partagée disponible par thread (si vous envisagez de l'utiliser indépendamment dans tous les threads).

Une particularité de la mémoire partagée est qu'elle est adressée de la même manière pour toutes les tâches d'un bloc (Fig. 7). Il s'ensuit qu'il peut être utilisé pour échanger des données entre des threads d'un seul bloc.

Il est garanti que lors de l'exécution d'un bloc sur un multiprocesseur, le contenu de la mémoire partagée sera préservé. Cependant, après la modification d'un bloc sur un multiprocesseur, il n'est pas garanti que le contenu de l'ancien bloc soit conservé. Par conséquent, vous ne devriez pas essayer de synchroniser les tâches entre les blocs, en laissant des données dans la mémoire partagée et en espérant leur sécurité.

Les variables déclarées avec le qualificatif __shared__ sont allouées en mémoire partagée.

Shared__ float mem_shared;

Il convient de souligner encore une fois qu'il n'existe qu'une seule mémoire partagée par bloc. Par conséquent, si vous devez l'utiliser simplement comme cache géré, vous devez accéder à différents éléments du tableau, par exemple, comme ceci :

float x = mem_shared;

Où threadIdx.x est l'index x du thread à l'intérieur du bloc.

3.2.4. Mémoire constante

La mémoire constante est mise en cache, comme le montre la Fig. 4. Le cache existe en une seule copie pour un multiprocesseur, ce qui signifie qu'il est commun à toutes les tâches du bloc. Sur l'hôte, vous pouvez écrire quelque chose dans la mémoire constante en appelant la fonction cudaMemcpyToSymbol. Depuis l'appareil, la mémoire constante est en lecture seule.

La mémoire constante est très pratique à utiliser. Vous pouvez y placer des données de n'importe quel type et les lire à l'aide d'une simple affectation.

#définir N 100

Constant__ int gpu_buffer[N];

vider la fonction hôte()

int cpu_buffer[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int)*N);

// __global__ signifie que device_kernel est un noyau qui peut être exécuté sur le GPU

Global__ void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer ;

// gpu_buffer = a; ERREUR! la mémoire constante est en lecture seule

La mémoire constante utilisant un cache, son accès est généralement assez rapide. Le seul mais très gros inconvénient de la mémoire constante est que sa taille n'est que de 64 Ko (pour l'ensemble de l'appareil). Cela suggère qu’il est logique de stocker seulement une petite quantité de données fréquemment utilisées dans la mémoire contextuelle.

3.2.5. Mémoire de textures

La mémoire de texture est mise en cache (Fig. 4). Il n'existe qu'un seul cache pour chaque multiprocesseur, ce qui signifie que ce cache est commun à toutes les tâches du bloc.

Le nom de mémoire de texture (et, malheureusement, de fonctionnalité) est hérité des concepts de « texture » et de « texturing ». La texturation est le processus d'application d'une texture (juste une image) à un polygone pendant le processus de rastérisation. La mémoire de texture est optimisée pour l'échantillonnage de données 2D et possède les capacités suivantes :

récupération rapide de valeurs de taille fixe (octet, mot, double ou quadruple mot) à partir d'un tableau unidimensionnel ou bidimensionnel ;

adressage normalisé avec des nombres flottants dans l'intervalle. Vous pouvez ensuite les sélectionner à l'aide d'un adressage normalisé. La valeur résultante sera un mot de type float4 mappé à l'intervalle ;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); // allouer de la mémoire dans le GPU

// définition des paramètres de texture texture

Texture.addressMode = cudaAddressModeWrap; // mode Envelopper

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //valeur la plus proche

Texture.normalized = faux ; // n'utilise pas d'adressage normalisé

CudaBindTexture (0, texture, gpu_memory, N ) // désormais cette mémoire sera considérée comme une mémoire de texture


CudaMemcpy (gpu_memory, cpu_buffer, N*sizeof (uint 4), cudaMemcpyHostToDevice ); // copie les données dansGPU

// __global__ signifie que device_kernel est le noyau à paralléliser

Global__ void device_kernel()

uint4 a = tex1Dfetch(texture,0); // vous ne pouvez sélectionner les données que de cette façon !

uint4 b = tex1Dfetch(texture,1);

int c = a.x * par;

...

3.3. Exemple simple

À titre d'exemple simple, considérons le programme cppIntegration du SDK CUDA. Il démontre des techniques pour travailler avec CUDA, ainsi que l'utilisation de nvcc (un sous-ensemble spécial du compilateur C++ de Nvidia) en combinaison avec MS Visual Studio, ce qui simplifie grandement le développement de programmes sur CUDA.

4.1. Décomposez correctement votre tâche

Toutes les tâches ne sont pas adaptées aux architectures SIMD. Si votre tâche ne s'y prête pas, cela ne vaut peut-être pas la peine d'utiliser un GPU. Mais si vous êtes déterminé à utiliser un GPU, vous devriez essayer de décomposer l'algorithme en morceaux pouvant être exécutés efficacement dans le style SIMD. Si nécessaire, modifiez l'algorithme pour résoudre votre problème, proposez-en un nouveau - un qui conviendrait bien à SIMD. Un exemple de domaine approprié pour l'utilisation du GPU est la mise en œuvre de l'addition pyramidale d'éléments de tableau.

4.2. Sélection du type de mémoire

Placez vos données dans une texture ou une mémoire constante si toutes les tâches d'un même bloc accèdent au même emplacement mémoire ou à des zones rapprochées. Les données 2D peuvent être traitées efficacement à l'aide des fonctions text2Dfetch et text2D. La mémoire de texture est spécifiquement optimisée pour l'échantillonnage 2D.

Utilisez la mémoire globale en combinaison avec la mémoire partagée si toutes les tâches accèdent de manière aléatoire à des emplacements de mémoire différents et largement séparés (avec des adresses ou des coordonnées très différentes s'il s'agit de données 2D/3D).

mémoire globale => mémoire partagée

Fils de synchronisation();

Traiter les données dans la mémoire partagée

Fils de synchronisation();

mémoire globale<= разделяемая память

4.3. Activer les compteurs de mémoire

L'indicateur du compilateur --ptxas-options=-v vous permet d'indiquer exactement combien et quel type de mémoire (registres, partagée, locale, constante) vous utilisez. Si le compilateur utilise la mémoire locale, vous le savez certainement. L'analyse des données sur la quantité et le type de mémoire utilisée peut grandement vous aider à optimiser votre programme.

4.4. Essayez de minimiser l'utilisation des registres et de la mémoire partagée

Plus le noyau utilise de registres ou de mémoire partagée, moins de threads (ou plutôt de warps) peuvent être exécutés simultanément sur un multiprocesseur, car Les ressources multiprocesseurs sont limitées. Par conséquent, une légère augmentation de l'occupation des registres ou de la mémoire partagée peut dans certains cas entraîner une baisse des performances de moitié - précisément parce que désormais exactement la moitié moins de warps sont exécutés simultanément sur un multiprocesseur.

4.5. Mémoire partagée au lieu de locale

Si le compilateur Nvidia, pour une raison quelconque, a alloué des données à la mémoire locale (cela est généralement perceptible par une très forte baisse des performances dans des endroits où il n'y a rien de gourmand en ressources), découvrez exactement quelles données se sont retrouvées dans la mémoire locale et placez-les dans partagé mémoire).

Souvent, le compilateur place une variable dans la mémoire locale si elle n'est pas fréquemment utilisée. Par exemple, il s’agit d’une sorte d’accumulateur dans lequel vous accumulez une valeur en calculant quelque chose dans une boucle. Si la boucle est volumineuse en termes de volume de code (mais pas en temps d'exécution !), alors le compilateur peut placer votre accumulateur dans la mémoire locale, car il est relativement rarement utilisé et les registres sont peu nombreux. Dans ce cas, la perte de performances peut être perceptible.

Si vous utilisez très rarement une variable, il est préférable de la placer explicitement dans la mémoire globale.

Bien qu'il puisse sembler pratique pour le compilateur de placer automatiquement de telles variables dans la mémoire locale, ce n'est en réalité pas le cas. Il sera difficile de trouver le goulot d'étranglement lors des modifications ultérieures du programme si la variable commence à être utilisée plus souvent. Le compilateur peut ou non transférer une telle variable vers la mémoire d'enregistrement. Si le modificateur __global__ est spécifié explicitement, le programmeur est plus susceptible d'y prêter attention.

4.6. Dérouler des boucles

Le déroulement de la boucle est une technique de performance standard dans de nombreux systèmes. Son essence est d'effectuer plus d'actions à chaque itération, réduisant ainsi le nombre total d'itérations, et donc le nombre de branches conditionnelles que le processeur devra effectuer.

Voici comment dérouler la boucle pour trouver la somme d'un tableau (par exemple, un entier) :

int un[N]; somme entière ;

pour (int je = 0; je

Bien entendu, les boucles peuvent également être déroulées manuellement (comme indiqué ci-dessus), mais il s'agit d'un travail improductif. Il est préférable d'utiliser des modèles C++ en combinaison avec des fonctions en ligne.

modèle

classe ArraySumm

Device__ statique T exec(const T* arr) ( return arr + ArraySumm (arr+1); )

modèle

classe ArraySumm<0,T>

Device__ static T exec(const T* arr) ( return 0; )

pour (int je = 0; je

somme+= TableauSumm<4,int>::exec(a);

Une fonctionnalité intéressante du compilateur nvcc est à noter. Le compilateur intégrera toujours les fonctions de type __device__ par défaut (il existe une directive spéciale __noinline__ pour remplacer cela).

Par conséquent, vous pouvez être sûr qu'un exemple comme celui ci-dessus se déroulera en une simple séquence d'instructions et ne sera en aucun cas inférieur en efficacité au code écrit à la main. Cependant, dans le cas général (pas nvcc), vous ne pouvez pas en être sûr, car inline n'est qu'une indication adressée au compilateur qu'il peut ignorer. Par conséquent, il n’est pas garanti que vos fonctions seront intégrées.

4.7. Alignement des données et échantillonnage de 16 octets

Alignez les structures de données sur des limites de 16 octets. Dans ce cas, le compilateur pourra utiliser des instructions spéciales qui chargeront les données 16 octets à la fois.

Si la structure fait 8 octets ou moins, vous pouvez l'aligner sur 8 octets. Mais dans ce cas, vous pouvez sélectionner deux variables à la fois en combinant deux variables de 8 octets dans une structure à l'aide d'une union ou d'un transfert de pointeur. Le transtypage doit être utilisé avec prudence car le compilateur peut placer les données dans la mémoire locale plutôt que dans des registres.

4.8. Conflits de banque de mémoire partagée

La mémoire partagée est organisée sous forme de 16 (au total !) banques de mémoire avec un pas de 4 octets. Lors de l'exécution d'un pool de threads warp sur un multiprocesseur, il est divisé en deux moitiés (si warp-size = 32) de 16 threads, qui accèdent tour à tour à la mémoire partagée.

Les tâches dans les différentes moitiés du Warp n'entrent pas en conflit sur la mémoire partagée. Du fait que les tâches de la moitié du pool de distorsion accéderont aux mêmes banques de mémoire, des collisions se produiront et, par conséquent, une diminution des performances. Les tâches situées dans la moitié de la chaîne peuvent accéder à différentes parties de la mémoire partagée avec une certaine étape.

Les pas optimaux sont 4, 12, 28, ..., 2^n-4 octets (Fig. 8).



Riz. 8. Étapes optimales.

Les étapes non optimales sont 1, 8, 16, 32, ..., 2^n octets (Fig. 9).



Riz. 9. Étapes sous-optimales

4.9. Minimiser les mouvements de données de l'hôte<=>Appareil

Essayez de transférer le moins possible les résultats intermédiaires vers l'hôte pour les traiter en utilisant le processeur. Implémentez, sinon l'intégralité de l'algorithme, du moins sa partie principale sur le GPU, laissant uniquement les tâches de contrôle au CPU.

5. Bibliothèque mathématique portable CPU/GPU

L'auteur de cet article a écrit une bibliothèque portable MGML_MATH pour travailler avec des objets spatiaux simples, dont le code est opérationnel à la fois sur l'appareil et sur l'hôte.

La bibliothèque MGML_MATH peut être utilisée comme cadre pour écrire des systèmes portables (ou hybrides) CPU/GPU pour calculer des problèmes physiques, graphiques ou autres problèmes spatiaux. Son principal avantage est que le même code peut être utilisé à la fois sur le CPU et sur le GPU, et en même temps, la vitesse est la principale exigence de la bibliothèque.

6 . Littérature

Chris Kaspersky. Techniques d'optimisation de programme. Utilisation efficace de la mémoire. - Saint-Pétersbourg : BHV-Pétersbourg, 2003. - 464 p. : ill.

Guide de programmation CUDA 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

Guide de programmation CUDA 1.1. pages 14-15

Guide de programmation CUDA 1.1. page 48

Selon la théorie de l'évolution de Darwin, le premier singe (si
pour être précis - homo antecessor, prédécesseur humain) devenu plus tard
en nous. Centres informatiques de plusieurs tonnes 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 sera pas inférieur en performances au premier. Les machines à écrire antédiluviennes sont devenues
à imprimer n'importe quoi et sur n'importe quoi (même sur le corps humain)
appareils multifonctionnels. Les géants des processeurs ont soudainement décidé de se murer
noyau graphique en "pierre". Et les cartes vidéo ont commencé non seulement à afficher une image avec
qualité FPS et graphique acceptable, mais effectue également toutes sortes de calculs. Oui
encore comment produire ! La technologie du calcul multithread utilisant les GPU sera discutée.

Pourquoi un GPU ?

Je me demande pourquoi ils ont décidé de transférer toute la puissance de calcul vers les graphiques
adaptateur? Comme vous pouvez le constater, les processeurs sont toujours à la mode et il est peu probable qu'ils abandonnent leur chaleur.
lieu. Mais le GPU a quelques atouts dans sa manche, ainsi qu'un joker et quelques manches.
assez. Un processeur central moderne est conçu pour atteindre un maximum
performances lors du traitement de données entières et à virgule flottante
virgule, sans se soucier particulièrement du traitement parallèle des informations. Au même
temps, l'architecture de la carte vidéo vous permet de « paralléliser » rapidement et sans problème
traitement de l'information. D'une part, les polygones sont calculés (grâce au convoyeur 3D),
d'autre part, le traitement de la texture des pixels. Il est clair qu’il existe une « harmonie
répartition de la charge dans le cœur de la carte. De plus, les performances de la mémoire et du processeur vidéo
plus optimale que la combinaison « RAM-cache-processeur ». Le moment où une unité de données
dans la carte vidéo commence à être traité par un processeur de flux GPU, un autre
l'unité est chargée en parallèle dans une autre et, en principe, il est facile de réaliser
Charge GPU comparable à la bande passante du bus,
cependant, pour que cela se produise, les convoyeurs doivent être chargés uniformément, sans
toutes les transitions et branches conditionnelles. Le processeur central, de par son
la polyvalence nécessite un cache complet pour ses besoins de traitement
information.

Les experts ont réfléchi au travail des GPU dans le calcul parallèle et
mathématiques et a proposé la théorie selon laquelle de nombreux calculs scientifiques sont similaires à bien des égards à
Traitement graphique 3D. De nombreux experts estiment que le facteur fondamental dans
développement GPGPU (Calcul à usage général sur GPU – universel
calculs à l'aide d'une carte vidéo) fut l'émergence du projet Brook GPU en 2003.

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

CUDA !

La prémonition des perspectives de développement a forcé DMLA Et Nvidia
accrochez-vous au GPU Brook comme un pitbull. Si nous omettons la politique de marketing, alors
En mettant tout 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 « seigneurs du FPS » se sont séparés face à la pierre d’achoppement, chacun à sa manière.
chemin, mais le principe de base est resté inchangé : faire des calculs
en utilisant le GPU. Et maintenant, nous allons examiner de plus près la technologie « verte » - CUDA
(Calculer une architecture de périphérique unifiée).

Le travail de notre « héroïne » consiste à fournir une API, deux à la fois.
Le premier est de haut niveau, CUDA Runtime, qui représente les fonctions qui
sont décomposés en niveaux plus simples et transmis à l'API inférieure - Pilote CUDA. Donc
que l’expression « haut niveau » est exagérée à appliquer au processus. Tout le sel est
exactement dans le pilote, et les bibliothèques aimablement créées vous aideront à l'obtenir
développeurs Nvidia: CUBLAS (outils de calculs mathématiques) et
FFT (calcul utilisant l'algorithme de Fourier). Bon, passons à la pratique
parties du matériau.

Terminologie CUDA

Nvidia fonctionne avec des définitions très uniques pour l'API CUDA. Ils
diffèrent des définitions utilisées pour travailler avec un processeur central.

Fil– un ensemble de données qui doivent être traitées (non
nécessite des ressources de traitement importantes).

Chaîne– un groupe de 32 threads. Les données sont traitées uniquement
les déformations, donc une déformation est la quantité minimale de données.

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

Grille est une collection de blocs. Cette division de données
utilisé uniquement pour améliorer les performances. Donc, si le numéro
Si le multiprocesseur est grand, les blocs seront exécutés en parallèle. Si avec
pas de chance avec la carte (les développeurs recommandent d'utiliser
adaptateur non inférieur à GeForce 8800 GTS 320 Mo), les blocs de données seront alors traités
séquentiellement.

NVIDIA introduit également des concepts tels que noyau, hôte
Et appareil.

Nous travaillons !

Pour travailler pleinement avec CUDA, vous avez besoin de :

1. Connaître la structure des cœurs de shader GPU, puisque l'essence de la programmation
consiste à 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é à plusieurs reprises "l'intérieur" de la carte vidéo
différemment de ce que nous avons l’habitude de voir. Alors, bon gré mal gré, il faudra tout étudier
subtilités de l'architecture. Regardons la structure de la légendaire « pierre » G80 GeForce 8800
GTX.

Le noyau du shader se compose de huit clusters TPC (Texture Processor Cluster).
processeurs de texture (donc, GeForce GTX 280– 15 cœurs, 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 (tous
16) se compose d'un frontal (résout les problèmes 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 (super function unit). Pour chaque battement (unité
temps), le frontal sélectionne la chaîne et la traite. Pour que toute la distorsion coule
(je vous le rappelle, il y en a 32) traités, 32/8 = 4 cycles sont nécessaires en fin de convoyeur.

Chaque multiprocesseur possède ce qu'on appelle une mémoire partagée.
Sa taille est de 16 kilo-octets et offre au programmeur une totale liberté
Actions. Distribuez comme vous le souhaitez :). La mémoire partagée assure la communication entre les threads
un bloc et n’est pas destiné à fonctionner avec des pixel shaders.

Les SM peuvent également accéder au GDDR. Pour ce faire, ils ont reçu 8 kilo-octets chacun.
mémoire cache qui stocke toutes les choses les plus importantes pour le travail (par exemple, l'informatique
constantes).

Le multiprocesseur dispose de 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. Il en ressort clairement que G80
peut traiter un maximum de 32*16*24 = 12288 threads par unité de temps. Tu ne peux pas t'empêcher
prendre en compte ces chiffres lors de l'optimisation future du programme (à une échelle
– la taille du bloc, d’autre part – le nombre de threads). L'équilibre des paramètres peut jouer un rôle
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. Prise en compte de toutes les subtilités de la structure de la carte vidéo GPU plus
bonnes compétences en programmation, vous pouvez créer des
outil de calcul parallèle. Au fait, à propos de la programmation...

La programmation

Pour la « créativité » avec CUDA, vous avez besoin Carte vidéo GeForce pas inférieure
épisode huit. AVEC

site officiel, vous devez télécharger trois packages logiciels : pilote de
Prise en charge de CUDA (chaque système d'exploitation a le sien), le package SDK CUDA lui-même (le deuxième
version 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
j'ai choisi Vista Ultimate Edition x64 (pour l'avenir, je dirai que le système s'est comporté
Tout simplement génial). Au moment d'écrire ces lignes, c'était pertinent pour le travail
Pilote ForceWare 177.35. Utilisé comme un ensemble d'outils
Le progiciel Borland C++ 6 Builder (bien que tout environnement fonctionnant avec
langue C).

Il sera facile pour une personne connaissant la langue de s'habituer à un nouvel environnement. Tout ce qui est requis est
rappelez-vous les paramètres de base. Mot clé _global_ (placé avant la fonction)
indique que la fonction appartient au noyau. Elle sera appelée 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 du maillage, la taille du bloc et quel sera le noyau
appliqué. Par exemple, la ligne _global_ void saxpy_parallel<<>>, où X –
la taille de la grille, et Y est la taille du bloc, spécifie ces paramètres.

Le symbole _device_ signifie que la fonction sera appelée par le cœur graphique, également appelé
suivra toutes les instructions. Cette fonction est située dans la mémoire du multiprocesseur,
il est donc impossible d'obtenir son adresse. Le préfixe _host_ signifie que l'appel
et le traitement aura lieu uniquement avec la participation du CPU. Il faut tenir compte du fait que _global_ et
Les _appareils_ ne peuvent pas s'appeler ni s'appeler eux-mêmes.

De plus, le langage pour CUDA dispose d'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. D'abord, il est pris
code destiné exclusivement au processeur central et soumis à
compilation standard et autre code destiné à l'adaptateur graphique
réécrit dans le langage intermédiaire PTX (un peu comme l'assembleur) pour
identifier les erreurs possibles. Après toutes ces « danses » la finale
traduction (traduction) des commandes dans un langage compréhensible pour GPU/CPU.

Trousse d'étude

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

Le navigateur CUDA SDK a été développé spécialement pour les débutants. N'importe qui peut
ressentez la puissance du calcul parallèle (le meilleur test pour
stabilité – les exemples fonctionnent sans artefacts ni plantages). L'application a
un grand nombre de mini-programmes indicatifs (61 « tests »). Pour chaque expérience il y a
Documentation détaillée du code du programme ainsi que des fichiers PDF. Il est immédiatement évident que les gens
ceux présents avec 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
données. Par exemple, analyser des tableaux multidimensionnels avec une carte vidéo GeForce 8800
GT Produit 512 Mo avec un bloc de 256 threads en 0,17109 millisecondes.
La technologie ne reconnaît pas les tandems SLI, donc si vous avez un duo ou un trio,
désactivez la fonction « appairage » avant de travailler, sinon CUDA n'en verra qu'un
appareil Double cœur AMD Athlon 64 X2(fréquence centrale 3000 MHz) même expérience
passe en 2,761528 millisecondes. Il s'avère que G92 est plus de 16 fois
plus vite qu'un rocher DMLA! Comme vous pouvez le constater, ce système est loin d’être extrême.
en tandem avec un système d'exploitation mal aimé des masses, cela fait du bien
résultats.

En plus du navigateur, il existe un certain nombre de programmes utiles à la société. Adobe
adapté ses produits aux nouvelles technologies. Maintenant Photoshop CS4 est complet
utilise le moins les ressources des adaptateurs graphiques (vous devez télécharger un fichier spécial
brancher). Avec des programmes tels que Badaboom Media Converter et RapiHD, vous pouvez
décoder la vidéo au format MPEG-2. Bon pour le traitement audio
L'utilitaire gratuit Accelero fera l'affaire. La quantité de logiciels adaptés à l'API CUDA,
va sans aucun doute croître.

Et à ce moment-là...

En attendant, vous lisez ce document, travailleurs acharnés des préoccupations des processeurs
développent leurs propres technologies pour intégrer les GPU dans les CPU. De l'exterieur DMLA Tous
c'est clair : ils ont une formidable expérience acquise avec ATI.

La création de « microdispositifs », Fusion, sera composée de plusieurs cœurs sous
nom de code Bulldozer et puce vidéo RV710 (Kong). Leur relation sera
réalisé grâce au bus HyperTransport amélioré. En fonction de la
nombre de cœurs et leurs caractéristiques de fréquence AMD prévoit de créer un prix global
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'application de la technologie
dans les appareils portables sera le premier défi pour les Canadiens. Avec développement
informatique parallèle, l’utilisation de telles « pierres » devrait être très populaire.

Intel un peu en retard avec son Larrabee. Des produits DMLA,
si rien ne se passe, ils apparaîtront dans les rayons des magasins fin 2009 - début
2010. Et la décision de l’ennemi ne sera révélée que dans presque deux
de l'année.

Larrabee aura un grand nombre (lire : des centaines) de cœurs. Au début
Il y aura également des produits conçus pour 8 à 64 cœurs. Ils sont très similaires aux Pentium, mais
assez fortement retravaillé. Chaque cœur dispose de 256 kilo-octets de cache L2
(sa taille augmentera avec le temps). La relation se fera à travers
Bus en anneau bidirectionnel 1024 bits. Intel dit que leur « enfant » sera
fonctionne parfaitement avec DirectX et Open GL API (pour les développeurs Apple), donc non
aucune intervention logicielle n'est requise.

Pourquoi je t'ai dit tout ça ? Il est évident que Larrabee et Fusion ne remplaceront pas
processeurs réguliers et stationnaires du marché, tout comme ils ne seront pas expulsés du marché
cartes vidéo. Pour les gamers et amateurs de sports extrêmes, le rêve ultime restera toujours
CPU multicœur et un tandem de plusieurs VGA haut de gamme. Mais quoi même
les fabricants de processeurs se tournent vers le calcul parallèle basé sur les principes
similaire à GPGPU, en dit long. En particulier, sur ce que tel
une technologie comme CUDA a le droit d’exister et, très probablement,
très populaire.

Un bref résumé

L'informatique parallèle utilisant une carte vidéo n'est qu'un bon outil
entre les mains d'un programmeur assidu. Les processeurs à peine 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 du grand public (on peut en dire autant de l'idée originale ATI/AMD).
À quoi cela ressemblera, l’avenir le dira. Donc 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 officiel et site Web de NVIDIA
GPGPU.com. Tous
les informations fournies sont en anglais, mais au moins merci qu'elles ne soient pas en anglais
Chinois Alors foncez ! J'espère que l'auteur vous a aidé au moins un peu
des voyages passionnants dans l'exploration de CUDA !




Partagez cet article:
Articles similaires

17 avril 2015
Comment ajouter de la musique dans GTA San Andreas ?

17 avril 2015
Que faire si GTA IV ne démarre pas

17 avril 2015
Instagram en chiffres - données utiles, faits et astuces

17 avril 2015
Changer le modem pour qu'il soit détecté comme un smartphone