M\u00eame les architectures optimis\u00e9es pour les hautes fr\u00e9quences comme le Prescott se sont cass\u00e9es les dents sur ce probl\u00e8me, et pour cause : cette fois il ne s\u2019agissait pas d\u2019un simple d\u00e9fi industriel, les fondeurs venaient tout simplement de se heurter aux lois de la physique. Certains Cassandres se sont mis alors \u00e0 proph\u00e9tiser la fin de la loi de Moore mais c\u2019\u00e9tait loin d\u2019\u00eatre le cas. Bien qu\u2019elle ait souvent \u00e9t\u00e9 d\u00e9tourn\u00e9e de son sens initial, le v\u00e9ritable sujet de la loi de Moore concerne le nombre de transistors sur une surface de silicium donn\u00e9e. Pendant longtemps la croissance du nombre de transistors des CPU s\u2019est certes accompagn\u00e9e d\u2019une augmentation de performance dans le m\u00eame temps, ce qui explique sans doute la confusion. Mais d\u00e9sormais les choses allaient se montrer plus compliqu\u00e9es : les architectes qui concevaient les CPU se heurtaient \u00e0 la loi des rendements d\u00e9croissants. Le nombre de transistors \u00e0 ajouter pour obtenir un gain de performance donn\u00e9 devenait de plus en plus important et menait tout droit dans une impasse.<\/p>\n
<\/span><\/span><\/p>\n \nPendant que les fabricants de CPU se creusaient la t\u00eate pour trouver une solution \u00e0 leurs probl\u00e8mes, les fabricants de GPU continuaient \u00e0 b\u00e9n\u00e9ficier plus que jamais des avantages de la loi de Moore.<\/p>\n <\/span><\/p>\n Pourquoi n\u2019\u00e9taient ils pas handicap\u00e9s comme leurs confr\u00e8res qui concevaient des CPU ? Pour une raison toute simple : les CPU sont con\u00e7us pour tirer le maximum de performances d\u2019un flux d\u2019instructions, celui-ci op\u00e8re sur des donn\u00e9es diverses (entiers, flottants), effectue des acc\u00e8s m\u00e9moire al\u00e9atoires, des branchements\u2026 Jusqu\u2019ici les architectes cherchaient \u00e0 extraire d\u2019avantage de parall\u00e9lisme d\u2019instructions, c\u2019est-\u00e0-dire \u00e0 lancer le plus d\u2019instructions possibles en parall\u00e8le. Ainsi le Pentium a introduit l\u2019ex\u00e9cution superscalaire en permettant de lancer, sous certaines conditions, deux instructions enti\u00e8res par cycle. Le Pentium Pro a pour sa part apport\u00e9 l\u2019ex\u00e9cution des instructions dans le d\u00e9sordre afin d\u2019utiliser au mieux les unit\u00e9s d\u2019ex\u00e9cution. Le probl\u00e8me est qu\u2019il ya une limite au parall\u00e9lisme qu\u2019il est possible d\u2019extraire d\u2019un flux s\u00e9quentiel d\u2019instructions. Par cons\u00e9quent, augmenter aveugl\u00e9ment le nombre d\u2019unit\u00e9s d\u2019ex\u00e9cution est vain car elles resteront pour la plupart inutilis\u00e9es la majeure partie du temps.<\/p>\n A l\u2019inverse le fonctionnement d\u2019un GPU est on ne peut plus simple : le travail consiste \u00e0 prendre un ensemble de polygones d\u2019un c\u00f4t\u00e9 et \u00e0 g\u00e9n\u00e9rer un ensemble de pixels de l\u2019autre. Les polygones et les pixels sont ind\u00e9pendants les uns des autres et peuvent donc \u00eatre trait\u00e9 par des unit\u00e9s parall\u00e8les. Un GPU peut donc se permettre de consacrer une grosse quantit\u00e9 de son die \u00e0 des unit\u00e9s de calcul qui, \u00e0 l\u2019inverse de celles d\u2019un CPU seront effectivement utilis\u00e9es.<\/p>\n <\/span><\/span><\/p>\n Autre point de divergence entre les deux unit\u00e9s : les acc\u00e8s m\u00e9moire d\u2019un GPU sont extr\u00eamement coh\u00e9rents : lorsqu\u2019un texel est lu, quelques cycles plus tard on lira le texel voisin, de la m\u00eame fa\u00e7on lorsqu\u2019un pixel est \u00e9crit quelques cycles plus tard un pixel voisin sera \u00e9crit. En organisant la m\u00e9moire de fa\u00e7on intelligente les performances se rapprochent fortement de la bande passante th\u00e9orique. Un GPU \u00e0 l\u2019inverse d\u2019un CPU n\u2019a donc pas besoin d\u2019un \u00e9norme cache, son r\u00f4le est principalement destin\u00e9 \u00e0 acc\u00e9l\u00e9rer les op\u00e9rations de texturing : quelques Ko sont donc suffisants pour contenir les quelques texels utilis\u00e9s dans les filtres bilin\u00e9aire ou trilin\u00e9aire.<\/p>\n <\/span><\/span><\/p>\n Ces deux mondes sont donc rest\u00e9s \u00e9trangers l\u2019un \u00e0 l\u2019autre pendant longtemps : on travaillait avec un (ou plusieurs) CPU, le GPU n\u2019\u00e9tait bon qu\u2019\u00e0 produire de jolies images rapidement. Mais un \u00e9v\u00e8nement va venir bouleverser tout \u00e7a : l\u2019apparition de la programmabilit\u00e9 dans les GPU. Encore une fois initialement il n\u2019y a pas lieu de s\u2019inqui\u00e9ter pour les CPU : les premiers GPU vant\u00e9s comme programmables (NV20, R200) sont loin d\u2019\u00eatre une menace : le nombre d\u2019instructions pour un programme reste limit\u00e9 \u00e0 une dizaine et ils travaillent sur des types de donn\u00e9es exotiques : 9 ou 12 bits \u00e0 virgule fixe.<\/p>\n <\/span><\/span><\/p>\n Mais la loi de Moore va encore faire son \u0153uvre : non seulement l\u2019augmentation du nombre de transistors permet d\u2019augmenter le nombre d\u2019unit\u00e9s de calcul mais elle permet \u00e9galement d\u2019augmenter leur flexibilit\u00e9. L\u2019apparition du NV30 va donc \u00eatre marquante \u00e0 plusieurs \u00e9gards. S\u2019il s\u2019agit d\u2019un GPU qui ne restera pas dans les annales pour les joueurs, il va apporter deux \u00e9l\u00e9ments importants pour commencer \u00e0 consid\u00e9rer le GPU autrement que comme un b\u00eate acc\u00e9l\u00e9rateur graphique :<\/p>\n A partir de ce moment toutes les conditions \u00e9taient r\u00e9unies pour attirer quelques chercheurs curieux et toujours \u00e0 la recherche de davantage de puissance de calcul.<\/p>\n L\u2019id\u00e9e d\u2019utiliser les acc\u00e9l\u00e9rateurs graphiques pour des calculs math\u00e9matiques n\u2019est pas r\u00e9cente. Il faut remonter aux ann\u00e9es 90 pour en trouver les premi\u00e8res traces. Initialement cela reste tr\u00e8s primitif, il s\u2019agit surtout d\u2019utiliser certaines fonctions c\u00e2bl\u00e9es du hardware comme le rasterizer ou le ZBuffer pour acc\u00e9l\u00e9rer des t\u00e2ches comme le path finding ou le trac\u00e9 de diagramme de Vorono\u00ef :<\/p>\n <\/span><\/p>\n <\/span><\/span>En 2003 avec l\u2019apparition de shaders \u00e9volu\u00e9s une \u00e9tape est franchie, cette fois il s\u2019agit d\u2019effectuer des calculs matriciels sur le hardware de l\u2019\u00e9poque. D\u00e8s cette ann\u00e9e l\u00e0 toute une section du Siggraph (\u00ab Computations on GPU<\/i> \u00bb) est d\u00e9di\u00e9e \u00e0 cette nouvelle frange de l\u2019informatique. Il ne s\u2019agit encore que des pr\u00e9misses de ce qui sera bient\u00f4t d\u00e9nomm\u00e9 GPGPU. Un premier tournant dans ce domaine sera l\u2019apparition de BrookGPU.<\/p>\n Pour bien comprendre le r\u00f4le de Brook il faut voir comment les choses se passaient avant son apparition : le seul moyen pour avoir acc\u00e8s aux ressources du GPU en 2003 \u00e9tait de passer par une des deux API graphiques : Direct3D ou OpenGL. Par cons\u00e9quent les chercheurs qui souhaitaient b\u00e9n\u00e9ficier de la puissance de calcul des GPU devaient travailler avec ces API. Le probl\u00e8me est que les personnes en question n\u2019\u00e9taient pas forc\u00e9ment expertes de la programmation graphique ce qui compliquait s\u00e9rieusement l\u2019acc\u00e8s \u00e0 cette technologie. L\u00e0 o\u00f9 un programmeur 3D parle de shader, de texture ou de fragment un adepte de la programmation parall\u00e8le parle de stream, de kernel, de scatter ou de gather. La premi\u00e8re difficult\u00e9 consiste donc \u00e0 trouver des analogies entre deux mondes distincts :<\/p>\n <\/span><\/p>\n \n<\/span>On le voit, m\u00eame avec ces analogies \u00e0 l\u2019esprit la t\u00e2che reste d\u00e9licate et c\u2019est l\u00e0 o\u00f9 Brook entre en jeu. Brook est un ensemble d\u2019extensions au langage C, \u00ab C with streams<\/i> \u00bb comme il fut pr\u00e9sent\u00e9 par ses cr\u00e9ateurs de l\u2019universit\u00e9 de Stanford. Concr\u00e8tement Brook propose d\u2019encapsuler toute la partie gestion de l\u2019API 3D pour exposer le GPU comme un coprocesseur de calculs parall\u00e8les. Pour cela Brook se compose de deux modules : d\u2019une part un compilateur qui \u00e0 partir d\u2019un fichier .br contenant du code C++ et des extensions, va g\u00e9n\u00e9rer du code C++ standard qui sera link\u00e9 \u00e0 un runtime offrant divers back-end (DirectX, OpenGL ARB, OpenGL NV3x, x86).<\/p>\n <\/span><\/span><\/p>\n Brook a eu plusieurs m\u00e9rites, le premier \u00e9tant de faire sortir le GPGPU de l\u2019ombre et de l\u2019exposer aux yeux du \u00ab grand public \u00bb. En effet \u00e0 l\u2019annonce de ce projet plusieurs sites web consacr\u00e9s \u00e0 l\u2019informatique se sont fait l\u2019\u00e9cho de l\u2019arriv\u00e9e de Brook, en simplifiant parfois de fa\u00e7on caricaturale la r\u00e9alit\u00e9 : \u00ab le CPU est mort, les GPU sont nettement plus puissants et bient\u00f4t pourront les remplacer \u00bb. 5 ans plus tard ce n\u2019est toujours pas le cas et soyons clairs : \u00e7a ne le sera jamais ! En revanche \u00e0 voir les \u00e9volutions successives des CPU qui s\u2019orientent vers de plus en plus de parall\u00e9lisme (toujours plus de cores, technologie de Simultaneous MultiThreading, \u00e9largissement des unit\u00e9s SIMD) et \u00e0 c\u00f4t\u00e9 de \u00e7a les GPU qui \u00e0 l\u2019inverse s\u2019orientent vers toujours plus de flexibilit\u00e9 (support des calculs flottants simple pr\u00e9cision, des calculs entiers et bient\u00f4t des calculs double pr\u00e9cision) il semble clair qu\u2019\u00e0 terme les deux sont destin\u00e9es \u00e0 se rencontrer. Qu\u2019adviendra-t-il alors ? Le GPU sera-t-il absorb\u00e9 par le CPU tout comme le coprocesseur arithm\u00e9tique avant lui ? C\u2019est possible. Intel et AMD travaillent sur des projets de ce genre mais le temps que cela arrive beaucoup de choses peuvent encore changer.<\/p>\n <\/span>Mais revenons \u00e0 notre sujet. Si le premier m\u00e9rite de Brook a \u00e9t\u00e9 de populariser la notion de GPGPU, cette API n\u2019a pas eu qu\u2019un r\u00f4le de communication : elle a aussi largement simplifi\u00e9 l\u2019acc\u00e8s aux ressources du GPU, permettant \u00e0 beaucoup plus de monde de commencer \u00e0 se former \u00e0 ce nouveau mod\u00e8le de programmation. En revanche malgr\u00e9 toutes les qualit\u00e9s de Brook il restait encore \u00e9norm\u00e9ment \u00e0 faire pour rendre le GPU cr\u00e9dible comme unit\u00e9 de calcul.<\/p>\n Un des probl\u00e8mes rencontr\u00e9s venait des diff\u00e9rentes couches d\u2019abstraction et en particulier de la surcharge de travail engendr\u00e9e par l\u2019API 3D qui pouvait \u00eatre non n\u00e9gligeable. Mais le v\u00e9ritable souci sur lequel les d\u00e9veloppeurs de Brook n\u2019avaient aucun contr\u00f4le venait de la compatibilit\u00e9. Il n\u2019est pas rare pour les fabricants de GPU d\u2019optimiser r\u00e9guli\u00e8rement leurs drivers surtout avec la forte concurrence qu\u2019ils s\u2019opposent. Si ces optimisations sont (la plupart du temps) b\u00e9n\u00e9fiques pour les joueurs, elles pouvaient en revanche briser du jour au lendemain la compatibilit\u00e9 de Brook. Difficile dans ces cas l\u00e0 d\u2019utiliser cette API dans du code de qualit\u00e9 industrielle, destin\u00e9 \u00e0 \u00eatre d\u00e9ploy\u00e9. Brook resta donc pendant longtemps l\u2019apanage des chercheurs et des programmeurs curieux.<\/p>\n Pour autant le succ\u00e8s d\u2019estime de Brook a suffit \u00e0 attirer l\u2019attention d\u2019ATI et de NVIDIA, les deux g\u00e9ants voyant dans cet int\u00e9r\u00eat naissant pour ce type d\u2019initiative l\u2019occasion d\u2019\u00e9largir un peu plus encore leur march\u00e9, touchant ainsi un nouveau secteur qui restait jusqu\u2019ici insensible \u00e0 leurs prouesses graphiques.<\/p>\n Certains chercheurs \u00e0 l\u2019origine de Brook ont donc rapidement rejoint les \u00e9quipes de d\u00e9veloppement de la firme de Santa Clara pour mettre sur pied une strat\u00e9gie globale visant \u00e0 cibler ce nouveau march\u00e9. L\u2019id\u00e9e consistait \u00e0 offrir un ensemble mat\u00e9riel\/logiciel adapt\u00e9 \u00e0 ce type de calculs : comme les d\u00e9veloppeurs de NVIDIA connaissent tous les secrets des GPU, plus question de reposer sur une API graphique, elle-m\u00eame ne communiquant avec le hardware que par le biais d\u2019un driver avec tous les probl\u00e8mes que cela implique comme nous l\u2019avons vu. L\u2019\u00e9quipe de d\u00e9veloppement de CUDA a donc d\u00e9velopp\u00e9 un ensemble de couches logicielles pour communiquer avec le GPU.<\/p>\n <\/span><\/span><\/p>\n Comme on le voit sur ce sch\u00e9ma CUDA offre deux API :<\/p>\n L\u2019API de haut niveau \u00e9tant impl\u00e9ment\u00e9e \u00ab au dessus \u00bb de l\u2019API bas niveau, chaque appel \u00e0 une fonction du runtime est d\u00e9compos\u00e9 en instructions plus basiques g\u00e9r\u00e9es par l\u2019API driver. Notons que ces deux API sont mutuellement exclusives : le programmeur doit utiliser l\u2019une ou l\u2019autre mais il est impossible de m\u00e9langer des appels de fonction de l\u2019une et de l\u2019autre. Lorsque l\u2019on parle d\u2019API de haut niveau il convient de relativiser : m\u00eame l\u2019API runtime reste ce que beaucoup consid\u00e9reraient comme d\u00e9j\u00e0 tr\u00e8s bas niveau, cependant elle offre des fonctions bien pratiques pour l\u2019initialisation ou la gestion des contextes. Malgr\u00e9 tout ne vous attendez pas \u00e0 beaucoup plus d\u2019abstraction : elle demande tout de m\u00eame une bonne connaissance des GPU NVIDIA et de la fa\u00e7on dont ils fonctionnent.<\/p>\n <\/span>L\u2019API driver est donc plus complexe \u00e0 g\u00e9rer, elle demande plus de travail pour lancer un traitement sur le GPU, mais en contrepartie elle est plus flexible, offrant un contr\u00f4le suppl\u00e9mentaire au programmeur qui le d\u00e9sire. Notons que les deux API sont capables de communiquer avec des ressources OpenGL ou Direct3D (9 seulement pour le moment). L\u2019utilit\u00e9 est \u00e9vidente : CUDA pourrait \u00eatre utilis\u00e9 pour g\u00e9n\u00e9rer des ressources (g\u00e9om\u00e9trie, textures proc\u00e9durales\u2026) qui seraient ensuite pass\u00e9es \u00e0 l\u2019API graphique ou \u00e0 l\u2019inverse on pourrait imaginer que l\u2019API 3D pourrait envoyer le r\u00e9sultat du rendu \u00e0 CUDA qui serait dans ce cas utilis\u00e9 pour effectuer un post traitement. Les exemples d\u2019interactions sont nombreux et l\u2019avantage est que les ressources restent stock\u00e9es dans la RAM du GPU sans n\u00e9cessiter de passer par le goulot d\u2019\u00e9tranglement du bus PCI-Express.<\/p>\n <\/span>A l\u2019inverse soulignons que le partage de ressources, en l\u2019occurrence la m\u00e9moire vid\u00e9o, avec les donn\u00e9es graphiques n\u2019est pas toujours idyllique et peut conduire \u00e0 quelques petits soucis : dans le cas d\u2019un changement de r\u00e9solution ou de profondeur de couleur, les donn\u00e9es graphiques ont la priorit\u00e9. Ainsi si les ressources pour le framebuffer doivent augmenter, le driver n\u2019h\u00e9sitera pas \u00e0 venir prendre celles allou\u00e9es aux applications utilisant CUDA, entra\u00eenant un plantage de ces derni\u00e8res. Pas tr\u00e8s \u00e9l\u00e9gant certes, mais il faut avouer que la situation devrait se pr\u00e9senter peu souvent. Puisque nous en sommes au chapitre des petits inconv\u00e9nients : l\u2019utilisation de plusieurs GPU par une application CUDA n\u00e9cessite de d\u00e9sactiver le mode SLI au pr\u00e9alable, sans cela un seul GPU sera visible au niveau de CUDA.<\/p>\n Enfin la troisi\u00e8me couche logicielle est un ensemble de biblioth\u00e8ques, deux pour \u00eatre pr\u00e9cis :<\/p>\n <\/span><\/p>\n \nAvant de nous plonger dans CUDA prenons soin de d\u00e9finir au pr\u00e9alable quelques termes qui pars\u00e8ment les documentations de NVIDIA. La firme Californienne a en effet choisi une terminologie bien particuli\u00e8re qui peut d\u00e9router. En premier lieu il faut d\u00e9finir ce qu\u2019est un thread<\/b> en CUDA car il n\u2019a pas tout \u00e0 fait le m\u00eame sens qu\u2019un thread CPU, et n\u2019est pas non plus \u00e9quivalent \u00e0 ce que nous appelons threads dans nos articles sur les GPU. Un thread sur le GPU consiste en un \u00e9l\u00e9ment de base des donn\u00e9es \u00e0 traiter. A l\u2019inverse des threads CPU, les threads CUDA sont extr\u00eamement \u00ab l\u00e9gers \u00bb ce qui signifie que le changement de contexte entre deux threads est une op\u00e9ration peu co\u00fbteuse.<\/p>\n Deuxi\u00e8me terme fr\u00e9quemment rencontr\u00e9 dans la documentation de CUDA : warp<\/b>. Cette fois pas de confusion ce terme n\u2019\u00e9voque rien si ce n\u2019est peut \u00eatre aux \u00ab Trekkies \u00bb ou aux adeptes de Warhammer. En r\u00e9alit\u00e9 pour la petite anecdote ce terme vient des machines \u00e0 tisser, il d\u00e9signe un ensemble de fils de cotons or (en anglais fil se dit\u2026 thread). Un warp en CUDA est donc un ensemble de 32 threads, il s\u2019agit de la taille minimale des donn\u00e9es trait\u00e9es de fa\u00e7on SIMD par un multiprocesseur en CUDA.<\/p>\n <\/span><\/span><\/p>\n Mais cette granularit\u00e9 n\u2019est toujours pas suffisante pour \u00eatre facilement utilisable par un programmeur, ainsi en CUDA on ne manipule pas directement des warps, on travaille avec des blocs<\/b> pouvant contenir de 64 \u00e0 512 threads.<\/p>\n Enfin ces blocs sont r\u00e9unis dans des grilles<\/b>. L\u2019int\u00e9r\u00eat de ce regroupement est que le nombre de blocs trait\u00e9s simultan\u00e9ment par le GPU est intimement li\u00e9 aux ressources du hardware comme nous le verrons plus loin. Le nombre de blocs dans une grille permet d\u2019abstraire totalement cette contrainte et d\u2019appliquer un kernel \u00e0 une grande quantit\u00e9 de threads en un seul appel, sans se soucier de ressources fix\u00e9es. Le runtime CUDA se charge de d\u00e9composer le tout pour nous. Ce mod\u00e8le est ainsi extr\u00eamement extensible : si un hardware a peu de ressources il ex\u00e9cute les blocs s\u00e9quentiellement, \u00e0 l\u2019inverse s\u2019il dispose d\u2019un tr\u00e8s grand nombre d\u2019unit\u00e9s il peut les traiter en parall\u00e8le. Le m\u00eame code permet donc de cibler \u00e0 la fois les GPU d\u2019entr\u00e9e de gamme, les GPU haut de gamme voire les GPU futurs.<\/p>\n <\/span><\/p>\n Les autres termes que vous rencontrerez fr\u00e9quemment dans l\u2019API CUDA sont utilis\u00e9s pour d\u00e9signer le CPU qui est ici appel\u00e9 host<\/b> (h\u00f4te) ou le GPU d\u00e9sign\u00e9 comme device<\/b> (p\u00e9riph\u00e9rique). Apr\u00e8s cette petite introduction qui, on l\u2019esp\u00e8re, ne vous aura pas trop refroidi, il est temps de passer aux choses s\u00e9rieuses !<\/p>\n \nFid\u00e8le lecteur de Tom\u2019s Hardware, l\u2019architecture des derniers GPU de NVIDIA n\u2019a plus aucun secret pour vous, si ce n\u2019est pas le cas courrez vite rattraper ce manque<\/a>. Avec CUDA, NVIDIA pr\u00e9sente son architecture d\u2019une fa\u00e7on l\u00e9g\u00e8rement diff\u00e9rente et expose certains d\u00e9tails qu\u2019il n\u2019\u00e9tait pas utile de d\u00e9voiler jusqu\u2019\u00e0 pr\u00e9sent.<\/p>\n <\/span><\/p>\n Comme vous pouvez le constater ci-dessus, le Shader Core de NVIDIA est compos\u00e9 de plusieurs clusters que NVIDIA nomme Texture Processor Cluster<\/b>. Une 8800GTX est compos\u00e9e par exemple de 8 clusters, une 8800GTS de 6 et ainsi de suite. Chaque cluster regroupe en fait une unit\u00e9 de texture et deux streaming multiprocessors<\/b>. Ces processeurs sont compos\u00e9s d\u2019un front-end de lecture\/d\u00e9codage et lancement des instructions et d\u2019un back-end compos\u00e9 d\u2019un ensemble de 8 unit\u00e9s de calcul et de 2 unit\u00e9s sp\u00e9ciales, au niveau desquels les instructions sont ex\u00e9cut\u00e9es de fa\u00e7on SIMD : la m\u00eame instruction est appliqu\u00e9e \u00e0 tous les threads du warp. NVIDIA baptise ce mode d\u2019ex\u00e9cution SIMT<\/b> pour Single Instruction Multiple Threads. Il est important de signaler que le back-end fonctionne \u00e0 une fr\u00e9quence double de celle du front-end.<\/p>\n En pratique la partie qui ex\u00e9cute les instructions appara\u00eet donc deux fois \u00ab plus large \u00bb qu\u2019elle ne l\u2019est (c\u2019est-\u00e0-dire comme une unit\u00e9 SIMD 16 voies au lieu de 8 voies). Le mode de fonctionnement des streaming multiprocessors est le suivant : \u00e0 chaque cycle un warp pr\u00eat \u00e0 \u00eatre ex\u00e9cut\u00e9 est s\u00e9lectionn\u00e9 par le front-end, qui lance l\u2019ex\u00e9cution d\u2019une instruction. Pour appliquer l\u2019instruction \u00e0 l\u2019ensemble des 32 threads du warp le back-end mettra quatre cycles mais comme il fonctionne \u00e0 une fr\u00e9quence double du front-end, il ne se sera ex\u00e9cut\u00e9 que deux cycles de son point de vue. Pour \u00e9viter que le front-end ne reste inutilis\u00e9 un cycle et maximiser l\u2019utilisation du hardware l\u2019id\u00e9al est donc d\u2019alterner les types d\u2019instructions tous les cycles : un cycle une instruction classique et l\u2019autre une instruction de type SFU.<\/p>\n Chaque multiprocesseur dispose \u00e9galement d\u2019un certain nombre de ressources qu\u2019il est utile de conna\u00eetre afin de les utiliser au mieux. Ainsi ils sont \u00e9quip\u00e9s d\u2019une petite zone m\u00e9moire appel\u00e9e Shared Memory<\/b> d\u2019une taille de 16 Ko par multiprocesseur. Cette m\u00e9moire n\u2019est pas une m\u00e9moire cache : sa gestion est enti\u00e8rement \u00e0 la charge du programmeur. En cela elle se rapproche de la Local Store des SPU du Cell. Cette sp\u00e9cificit\u00e9 est particuli\u00e8rement int\u00e9ressante et traduit le fait que CUDA est bien un ensemble de technologies logicielles et mat\u00e9rielles. En effet cette zone m\u00e9moire n\u2019est pas utilis\u00e9e dans le cas des pixels shaders, comme le pr\u00e9cise NVIDIA avec humour \u00ab nous n\u2019appr\u00e9cions pas que les pixels parlent les uns avec les autres \u00bb.<\/p>\n \nCette shared memory<\/i>offre un moyen aux threads d\u2019un m\u00eame bloc<\/u> de communiquer. Il est important de souligner la restriction : tous les threads d\u2019un m\u00eame bloc sont en effet garantis d\u2019\u00eatre ex\u00e9cut\u00e9s par le m\u00eame multiprocesseur. A l\u2019inverse l\u2019attribution des blocs aux diff\u00e9rents multiprocesseurs est compl\u00e8tement ind\u00e9finie, deux threads de blocs distincts ne peuvent donc pas communiquer durant leur ex\u00e9cution. Bien utiliser cette m\u00e9moire est donc compliqu\u00e9 mais peut se r\u00e9v\u00e9ler payant car, hormis le cas o\u00f9 plusieurs threads tentent d\u2019acc\u00e9der \u00e0 une m\u00eame banque m\u00e9moire ce qui provoque un conflit, le reste du temps l\u2019acc\u00e8s \u00e0 la shared memory<\/i> s\u2019av\u00e8re aussi performant que l\u2019acc\u00e8s aux registres.<\/p>\n <\/span><\/p>\n La shared memory<\/i> n\u2019est pas la seule m\u00e9moire auquel les multiprocesseurs ont acc\u00e8s, ils peuvent \u00e9videmment avoir recours \u00e0 la m\u00e9moire vid\u00e9o mais celle-ci offre une bande passante plus basse et une latence plus \u00e9lev\u00e9e. Par cons\u00e9quent pour limiter les acc\u00e8s trop fr\u00e9quents \u00e0 cette m\u00e9moire NVIDIA a donc dot\u00e9 ses multiprocesseurs de cache (d\u2019une taille d\u2019environ 8 Ko par multiprocesseur) pour l\u2019acc\u00e8s aux constantes ou aux textures.<\/p>\n <\/span><\/p>\n Les multiprocesseurs disposent aussi de 8192 registres \u00e0 partager entre tous les threads de tous les blocs actifs sur ce multiprocesseur. Le nombre de blocs actifs par multiprocesseur pour sa part ne peut pas d\u00e9passer 8, le nombre de warps actifs \u00e9tant pour sa part limit\u00e9 \u00e0 24 (768 threads). Une 8800 GTX peut donc avoir jusqu\u2019\u00e0 12 288 threads en cours de traitement \u00e0 chaque instant. Conna\u00eetre toutes ces limites peut sembler r\u00e9barbatif mais est utile afin de bien dimensionner son probl\u00e8me en fonction des ressources disponibles.<\/p>\n Optimiser un programme CUDA consiste donc essentiellement \u00e0 \u00e9quilibrer au mieux le nombre de blocs et leur taille : plus de threads par blocs s\u2019av\u00e8re utile pour mieux masquer la latence des op\u00e9rations m\u00e9moires mais d\u2019un autre c\u00f4t\u00e9 cela diminue le nombre de registres disponibles par threads. De plus un bloc de 512 threads serait particuli\u00e8rement peu efficace car seul un bloc pourrait \u00eatre actif sur un multiprocesseur, g\u00e2chant ainsi potentiellement 256 threads. NVIDIA conseille donc d\u2019utiliser des blocs de 128 \u00e0 256 threads qui offrent le meilleur compromis entre masquage de la latence et nombre de registres suffisant pour la plupart des kernels.<\/p>\n \nD\u2019un point de vue logiciel CUDA consiste en un ensemble d\u2019extensions au langage C, qui \u00e9voquent des souvenirs de BrookGPU, et en quelques appels d\u2019API sp\u00e9cifiques. Au niveau des extensions on trouve notamment des qualificateurs s\u2019appliquant aux fonctions et aux variables. Le mot cl\u00e9 principal \u00e0 retenir est __global__<\/b>. Plac\u00e9 devant une fonction il indique que celle-ci est un kernel c\u2019est-\u00e0-dire une fonction qui va \u00eatre appel\u00e9e par le CPU et ex\u00e9cut\u00e9e par le GPU. Le qualificateur __device__<\/b> pour sa part d\u00e9signe une fonction qui sera ex\u00e9cut\u00e9e par le GPU mais qui n\u2019est appelable que depuis le GPU (autrement dit depuis une autre fonction __device__ ou depuis une fonction __global__). Enfin le mot cl\u00e9 __host__<\/b> est optionnel, il d\u00e9signe une fonction qui est appel\u00e9e par le CPU et ex\u00e9cut\u00e9e sur le CPU, autrement dit une fonction traditionnelle.<\/p>\n Notons quelques restrictions associ\u00e9es aux fonctions __device__ ou __global__ : elles ne peuvent \u00eatre r\u00e9cursives (c\u2019est-\u00e0-dire s\u2019appeler elles m\u00eames) et elles ne peuvent pas avoir un nombre variable d\u2019arguments. Enfin les fonctions __device__ r\u00e9sidant dans l\u2019espace m\u00e9moire du GPU il est en toute logique impossible d\u2019obtenir leur adresse. Les variables disposent elles aussi de nouveaux qualificateurs permettant de contr\u00f4ler la zone m\u00e9moire dans laquelle elles seront stock\u00e9es. Ainsi une variable pr\u00e9c\u00e9d\u00e9e du mot cl\u00e9 __shared__<\/b> indique qu\u2019elle sera stock\u00e9e dans la shared memory des streaming multiprocessors.<\/p>\n L\u2019appel d\u2019une fonction __global__ est \u00e9galement un peu particulier. Il faut en effet d\u00e9finir lors de cet appel la configuration d\u2019ex\u00e9cution c\u2019est-\u00e0-dire plus concr\u00e8tement : la taille de la grille sur laquelle le kernel est appliqu\u00e9 et la taille de chaque bloc. Exemple un kernel dont la signature est la suivante :<\/p>\n __global__ void Func(float* parameter) ;<\/p>\n Sera appel\u00e9 ainsi :<\/p>\n Func<<>>(parameter) ;<\/p>\n Avec Dg comme dimension de grille et Db comme dimension d\u2019un bloc. Ces deux variables \u00e9tant d\u2019un nouveau type vectoriel introduit par CUDA.<\/p>\n L\u2019API CUDA offre quant \u00e0 elle essentiellement des fonctions de manipulation m\u00e9moire en VRAM : cudaMalloc pour allouer de la m\u00e9moire, cudaFree pour la lib\u00e9rer ou encore cudaMemcpy pour copier des donn\u00e9es entre RAM et VRAM et vice versa.<\/p>\n Terminons ce tour d\u2019horizon par la fa\u00e7on dont un programme CUDA est compil\u00e9, qui s\u2019av\u00e8re int\u00e9ressante. La compilation est effectu\u00e9e en plusieurs phases : tout d\u2019abord le code d\u00e9di\u00e9 au CPU est extrait du fichier et pass\u00e9 au compilateur standard. Le code d\u00e9di\u00e9 au GPU pour sa part est tout d\u2019abord converti en un langage interm\u00e9diaire : PTX. Ce langage interm\u00e9diaire est proche d\u2019un assembleur et permet donc d\u2019\u00e9tudier le code source g\u00e9n\u00e9r\u00e9 et de noter les inefficacit\u00e9s potentielles. Enfin la derni\u00e8re \u00e9tape traduit ce langage interm\u00e9diaire en commandes sp\u00e9cifiques au GPU et les encapsulent sous forme binaire dans l\u2019ex\u00e9cutable.<\/p>\n <\/span><\/span><\/p>\n \nDifficile, apr\u00e8s avoir ingurgit\u00e9 toute la documentation de NVIDIA, de r\u00e9sister \u00e0 la tentation de mettre les mains dans le cambouis. Apr\u00e8s tout quel meilleur moyen de juger une API que d\u2019essayer d\u2019\u00e9crire un petit programme l\u2019utilisant ? C\u2019est dans cette situation que se r\u00e9v\u00e8lent la plupart des probl\u00e8mes alors que tout semble parfait sur le papier. C\u2019est aussi le meilleur moyen de voir si nous avons bien assimil\u00e9 tous les concepts d\u00e9crits dans la documentation de CUDA.<\/p>\n Rien de plus facile actuellement que de se lancer dans un tel projet : on trouve des outils de tr\u00e8s bonne qualit\u00e9 gratuitement, pour ce test nous nous sommes donc bas\u00e9s sur Visual C++ Express 2005 qui offrait tout ce dont nous avions besoin. Le plus dur fut au final de trouver un programme suffisamment simple pour que nous puissions le porter sur le GPU sans y passer des semaines mais qui soit en m\u00eame temps suffisamment int\u00e9ressant pour que l\u2019op\u00e9ration ait un minimum d\u2019int\u00e9r\u00eat. Notre choix s\u2019est port\u00e9 finalement sur un bout de code dont nous disposions qui calculait \u00e0 partir d\u2019une heightmap, la normal map correspondante. Inutile de s\u2019\u00e9terniser sur les d\u00e9tails de la fonction qui ne pr\u00e9sentent pas particuli\u00e8rement d\u2019int\u00e9r\u00eat dans le cas pr\u00e9sent. Pour \u00eatre concis il suffit de dire qu\u2019il s\u2019agit d\u2019une convolution : pour chaque pixel de l\u2019image de d\u00e9part on applique une matrice qui va d\u00e9terminer, \u00e0 partir des pixels voisins selon une formule plus ou moins compliqu\u00e9e, la couleur du pixel r\u00e9sultant dans l\u2019image g\u00e9n\u00e9r\u00e9e. L\u2019avantage de cette fonction est qu\u2019elle est tr\u00e8s facilement parall\u00e9lisable, c\u2019est un cas id\u00e9al pour lequel CUDA pr\u00e9sente un int\u00e9r\u00eat. \n\n R\u00e9p\u00e9tons le encore une fois : l\u2019objectif de ce test \u00e9tait de se familiariser en pratique avec les outils du SDK CUDA, l\u2019id\u00e9e n\u2019\u00e9tait pas de faire un bench comparatif entre une version CPU et une version GPU. En temps que premier programme CUDA que nous allions \u00e9crire il ne fallait pas s\u2019attendre \u00e0 des merveilles d\u2019un point de vue performance. De plus, vu qu\u2019elle ne faisait pas partie d\u2019une section de code critique, la version CPU n\u2019\u00e9tait d\u00e9j\u00e0 pas optimis\u00e9e outre mesure, une comparaison directe des r\u00e9sultats n\u2019aurait pas vraiment d\u2019int\u00e9r\u00eat.<\/p>\n Malgr\u00e9 cela, nous avons quand m\u00eame choisi de mesurer le temps de calcul pour v\u00e9rifier si, malgr\u00e9 notre impl\u00e9mentation na\u00efve, il y avait un int\u00e9r\u00eat \u00e0 utiliser CUDA ou si le GPU ne s\u2019apprivoisait qu\u2019apr\u00e8s \u00e9norm\u00e9ment de pratique. La machine de test est notre machine de d\u00e9veloppement\u00a0: un ordinateur portable \u00e9quip\u00e9 d\u2019un Core 2 Duo T5450 et d\u2019une GeForce 8600M GT le tout tournant sous Vista. C\u2019est bien loin d\u2019une machine de guerre mais les r\u00e9sultats demeurent int\u00e9ressants car il s\u2019agit d\u2019un cas assez peu favorable pour le GPU\u00a0: il est bien pratique pour NVIDIA de montrer des acc\u00e9l\u00e9rations cons\u00e9quentes sur des syst\u00e8mes \u00e9quip\u00e9s de GPU monstrueux et disposant d\u2019une bande passante \u00e9norme, mais en pratique beaucoup des 70 millions de GPU CUDA \u00e9quipant des PC actuellement sont nettement moins puissants que \u00e7a, notre test nous place donc dans un cas pratique.<\/p>\n\n\n\n\n\n <\/p>\n\n Les r\u00e9sultats que nous avons obtenus sont les suivants pour le traitement d\u2019une image de 2048×2048\u00a0:<\/p>\n\n\n\n\n\n <\/p>\n\n Plusieurs observations sont \u00e0 extraire de ces r\u00e9sultats\u00a0: tout d\u2019abord vous noterez que nous avons \u00e9t\u00e9 m\u00e9disants car nous avons malgr\u00e9 tout modifi\u00e9 l\u2019impl\u00e9mentation initiale du CPU en la threadant. Comme nous l\u2019avons dit le code est id\u00e9al pour ce cas de figure, il suffit de d\u00e9composer l\u2019image initiale en autant de zones que de threads. Vous noterez que l\u2019on obtient une acc\u00e9l\u00e9ration quasiment lin\u00e9aire en passant de 1 \u00e0 2 threads sur notre CPU dual core ce qui traduit bien la nature fortement parall\u00e8le de ce programme de test. De fa\u00e7on assez inexpliqu\u00e9e la version 4 threads se r\u00e9v\u00e8le plus rapide alors que nous nous attendions au mieux \u00e0 ne voir aucune diff\u00e9rence sur notre processeur voire m\u00eame de fa\u00e7on plus logique \u00e0 une l\u00e9g\u00e8re perte d\u2019efficacit\u00e9 du fait du surcout engendr\u00e9 par la cr\u00e9ation des threads suppl\u00e9mentaires. Comment expliquer ce r\u00e9sultat\u00a0? Difficile \u00e0 dire, peut \u00eatre que l\u2019ordonnanceur de threads de Windows n\u2019est pas totalement innocent l\u00e0-dessous, en tout cas ce r\u00e9sultat \u00e9tait reproductible. Sur une texture aux dimensions plus r\u00e9duites (512×512) le gain obtenu en threadant est beaucoup moins sensible (35% environ au lieu de 100%) et le comportement de la version 4 threads est plus logique vu qu\u2019il n\u2019y a aucun gain par rapport \u00e0 la version 2 threads. Le GPU reste le plus rapide mais de fa\u00e7on moins sensible (la 8600M GT est 3 fois plus rapide que la version 2 threads).<\/p>\n\n\n\n\n\n <\/p>\n\nPendant ce temps\u2026<\/h2>\n
Vive le GeForce FX !<\/h2>\n
L\u2019apparition du GPGPU : pr\u00e9mices<\/h2>\n
BrookGPU<\/h2>\n
L\u2019API CUDA<\/h2>\n
Quelques d\u00e9finitions<\/h2>\n
CUDA d\u2019un point de vue mat\u00e9riel<\/h2>\n
Point de vue mat\u00e9riel (suite)<\/h2>\n
CUDA d\u2019un point de vue logiciel<\/h2>\n
En pratique<\/h2>\n
<\/span><\/td> <\/span><\/td><\/tr><\/table><\/div>\n<\/span>\nLe second avantage est que nous disposions d\u00e9j\u00e0 d\u2019une impl\u00e9mentation CPU \u00e0 laquelle nous pourrons facilement comparer le r\u00e9sultat de notre version CUDA, ce qui nous \u00e9vitait de r\u00e9inventer la roue selon la formule consacr\u00e9e de tout bon programmeur. En pratique lorsqu\u2019un programmeur vous dit \u00e7a c\u2019est surtout que le temps ainsi gagn\u00e9 pourra \u00eatre mis \u00e0 contribution de fa\u00e7on tr\u00e8s productive en jouant au jeu du moment ou en regardant la comp\u00e9tition sportive d\u2019actualit\u00e9\u2026<\/p>\n Performances<\/h2>\n
![\"Image](\"https:\/\/www.tomshardware.fr\/content\/uploads\/sites\/3\/2008\/06\/026-755x370.png\")
<\/a><\/span><\/span><\/span><\/span><\/p>\n\n\n\n\n\n