Intel concurrence NVIDIA avec son nouveau Xeon Phi Knights Landing

Annoncée il y a de cela un an, la nouvelle mouture des processeurs Intel Xeon Phi, nom de code Knights Landing, était annoncée : ces processeurs sont maintenant livrés aux fournisseurs de matériel HPC. Avec ces nouvelles puces, Intel vise le même segment que NVIDIA avec ses cartes graphiques (GPU) de génération Pascal : l’apprentissage profond, le calcul scientifique de haute performance (HPC).

Actuellement, quatre modèles sont disponibles, avec un nombre de cÅ“urs variable (de soixante-quatre à septante-deux) et des fréquences aussi variables (entre 1,3 et 1,5 GHz), des caractéristiques proches des GPU actuels comme NVIDIA Pascal. Tous sont livrés avec seize gigaoctets de mémoire vive (MCDRAM), empilée et très proche du processeur, pour une bande passante de presque cinq cents gigaoctets par seconde. Ces différentes puces représentent des compromis différents : la plus puissante (7290, 3,46 Tflops) est la plus chère (6294 $), avec un produit bien plus abordable au niveau du téraflops par euro (7210 à 2438 $ pour 2,66 Tflops) ; l’une optimise le ratio performance par watt (7250, 3,05 Tflops pour 215 W, là où le 7290 fournit 3,46 Tflops pour 245 W) et la dernière la mémoire disponible par cÅ“ur (7230, soixante-quatre cÅ“urs, comme le 7210, mais la mémoire a une fréquence plus élevée).

Au niveau des chiffres bruts, ces processeurs ne dépassent pas l’offre de NVIDIA : « Ã  peine » trois téraflops et demi pour le plus haut de gamme, quand les GPU NVIDIA Pascal atteignent cinq téraflops (et une bande passante de cinquante pour cent plus élevée, à sept cent vingt gigaoctets par seconde). Par contre, selon les applications, à cause de la différence d’architecture fondamentale, les résultats varient énormément (les Xeon Phi sont organisés comme des processeurs traditionnels : chaque cÅ“ur peut exécuter une instruction propre, contrairement aux GPU, où la même instruction est exécutée sur un grand nombre de cÅ“urs). Ainsi, pour de la simulation de dynamique moléculaire, pour le test de performance LAMPPS, un Xeon Phi de milieu de gamme (7250) a fonctionné cinq fois plus vite en consommant huit fois moins de mémoire qu’un GPU NVIDIA K80 (la génération précédant les Pascal). Pour de la visualisation par lancer de rayon, Intel indique être cinq fois plus rapide ; le facteur descend à trois pour de la simulation de risque financier. Ces comparaisons ne sont pas entièrement équitables, à cause de la différence d’âge entre les processeurs, mais donnent la tendance qu’Intel veut montrer. Ces résultats devront être confirmés par des indépendants pour être fiables.

Parmi les tests de performance, Intel s’est aussi orienté vers l’apprentissage profond, branche dans laquelle NVIDIA s’impose actuellement. Ce domaine est actuellement à la pointe de la recherche, avec des résultats de plus en plus intéressants : c’est grâce à des techniques de ce genre que Google a pu battre le joueur le plus fort au monde au jeu de go. Sur un même jeu de test, un Xeon Phi 7250 a obtenu sa réponse en cinquante fois moins de temps qu’un seul processeur traditionnel ; avec quatre tels Xeon Phi, les temps de calcul ont été réduits d’un facteur deux par rapport à quatre GPU NVIDIA K80.

Intel précise également qu’il est plus facile de programmer ses processeurs Xeon Phi : ils embarquent beaucoup plus de cÅ“urs, mais c’est la seule différence avec les processeurs habituels de nos PC, alors qu’il faut réécrire complètement son code (ou utiliser des API adaptées) pour les GPU, avec une phase d’optimisation du code qui nécessite des compétences plus spécifiques. La nouvelle génération de Xeon Phi apporte cependant une distinction plus importante : ces processeurs Intel pourront être utilisés comme processeurs principaux (pas seulement comme cartes d’extension), ce qui évite les opérations de transfert de données, très limitantes pour la performance des applications actuelles. Il reste cependant à déterminer si ces processeurs seront suffisamment rapides pour effectuer toutes les opérations des applications qui leur sont soumises (ils fonctionnent à une fréquence réduite par rapport aux processeurs habituels : ils excellent dans le traitement parallèle, mais pas en série, qui constitue parfois une part importante du code à exécuter).

Ni ces nouveaux Xeon Phi ni les GPU NVIDIA Pascal ne sont actuellement utilisés à grande échelle pour du calcul scientifique. Cependant, ces premiers résultats montrent que les deux jouent dans la même cour. Si les mesures d’Intel se généralisent, ils deviendront un concurrent plus que très sérieux de NVIDIA, notamment dans le marché en expansion de l’apprentissage profond ; s’ils ont en plus l’avantage du prix, la dominance de NVIDIA sera vite mise à mal.

Sources : Intel Takes on NVIDIA with Knights Landing Launch, Intel’s Knights Landing: Fresh x86 Xeon Phi lineup for HPC and AI, Intel Xeon Phi Knights Landing Now Shipping; Omni Path Update, Too.

Merci à Claude Leloup pour sa relecture orthographique.

NVIDIA annonce son architecture Pascal et CUDA 8

NVIDIA en parle depuis un certain temps et dégaine, peu après AMD, sa nouvelle générations de processeurs graphiques, dénommée Pascal, un nom qui fait référence à Blaise Pascal, mathématicien et physicien français du XVIIe siècle. Les chiffres bruts donnent déjà le tournis : un tel GPU est composé de dix-huit milliards de transistors (huit milliards pour la génération précédente, Maxwell).

Sur la technique, ce GPU (nommé GP100) est le premier gravé en 16 nm, avec la technologie FinFET+ de TSMC, les précédents l’étaient en 28 nm. Pour un processus aussi récent, la puce est énorme : 610 mm², à comparer au 600 mm² de Maxwell sur un processus bien maîtrisé. Cette surface est utilisée pour les 3840 cÅ“urs de calcul (3072 pour Maxwell), ayant accès à un total de 14 Mo de mémoire cache (6 Mo pour Maxwell). Les cÅ“urs ont aussi augmenté en fréquence : 1,5 GHz au lieu de 1,1 pour Maxwell. Au niveau de la mémoire principale, la technologie HBM2 est utilisée au lieu de la mémoire traditionnelle GDDR5. Une puce Pascal dispose de seize gigaoctets de mémoire (contre douze, il n’y a pas si longtemps, chez Maxwell, maintenant vingt-quatre).

NVIDIA peut donc en profiter pour intégrer bon nombre de nouvelles fonctionnalités, tournées vers le segment HPC, comme sa connectique NVLink. L’une des fonctionnalités très utiles pour monter en performance concerne les nombres à virgule flottante : contrairement à Maxwell, les unités de calcul de Pascal traitent nativement des nombres sur soixante-quatre bits et seize bits ; ainsi, une telle puce peut monter jusque vingt et un mille milliards d’opérations par seconde (21 TFlops) en seize bits, cinq mille milliards en soixante-quatre bits (avec un facteur d’à peu près quatre entre les deux : les unités sur soixante-quatre bits sont grosso modo composées de quatre unités sur seize bits).

Toujours dans le HPC, la mémoire unifiée, entre la mémoire principale de l’ordinateur et celle du processeur graphique, a aussi vu quelques améliorations. Elles se comptent au nombre de deux. Tout d’abord, l’espace d’adressage a été considérablement élargi, ce qui permet d’adresser toute la mémoire centrale et celle du GPU (et non seulement un espace équivalent à la mémoire du GPU). Ensuite, les GPU Pascal gèrent les défauts de page : si un noyau de calcul tente d’accéder à une page mémoire qui n’est pas encore chargée sur le GPU, elle le sera à ce moment-là, elle ne doit plus être copiée dès le lancement des calculs sur le GPU.

L’apprentissage profond, le cheval de bataille de NVIDIA

Toutes ces avancées sont jusque maintenant principalement prévues pour le marché du HPC, plus particulièrement de l’apprentissage de réseaux neuronaux profonds, à l’origine de miracles récents dans le domaine de l’intelligence artificielle.

NVIDIA avance un facteur d’accélération de douze dans un cas d’apprentissage d’un réseau neuronal avec sa nouvelle architecture par rapport à ses meilleurs résultats l’année dernière. Ce facteur est énorme en seulement une année, mais il faut remarquer que la comparaison proposée n’est pas toujours équitable : les nouveaux tests utilisent deux fois plus de GPU que l’année dernière, de nouvelles fonctionnalités ont fait leur apparition sur les GPU spécifiquement pour accélérer ce genre de calculs (les nombres à virgule flottante sur seize bits). Il n’empêche, le résultat reste impressionnant.

Ces avancées réunissent cinq « miracles », selon la terminologie de Jen-Hsun Huang :

  • la nouvelle architecture Pascal ;
  • la technologie de gravure de puces de TSMC ;
  • la nouvelle génération de puces mémoire empilées (HBM2) ;
  • la technologie d’interconnexion entre GPU NVLink ;
  • des améliorations algorithmiques pour tirer parti des progrès précédents.

NVIDIA DGX-1, un supercalculateur dans une boîte

Les tests de performance ont été effectués sur une machine assemblée par NVIDIA, dénommée DGX-1, un monstre de puissance. Elle couple deux processeurs Intel Xeon E5-2698 v3 et huit processeurs NVIDIA Pascal avec 512 Go de mémoire vive et 1,92 To de stockage en SSD. Globalement, la machine peut fournir 170 TFlops pour la modique somme de 129 000 $. NVIDIA la résume comme deux cent cinquante serveurs rassemblés dans un boîtier. Elle est d’ores et déjà en vente, mais les livraisons devraient débuter en juin. Des bibliothèques existantes ont été adaptées pour cette machine pour en tirer le meilleur parti, comme TensorFlow de Google.

CUDA 8 et autres avancées logicielles

Pascal vient avec des améliorations matérielles accessibles par CUDA 8, la nouvelle version de l’API de calcul sur GPU de NVIDIA. Par exemple, les opérations atomiques en mémoire, utiles pour les opérations en parallèle sur une même case mémoire sans risque de corruption, gèrent plus d’opérations sur les entiers (en parallèle avec une implémentation matérielle plus complète). Les développeurs peuvent exploiter une structure d’interconnexion NVLink explicitement depuis CUDA.

Le profilage d’applications est encore amélioré, notamment au niveau de la répartition de la charge entre le CPU et le GPU, pour déterminer les parties où l’optimisation serait la plus bénéfique globalement. Pour ce faire, le profileur visuel montre maintenant les dépendances entre les noyaux de calcul du GPU et les appels à l’API CUDA, afin de faciliter la recherche d’un chemin critique d’exécution.

De nouvelles bibliothèques font également leur apparition au sein de l’écosystème, comme nvGRAPH pour l’analyse de graphes : la planification du chemin d’un robot dans un graphe énorme (comme une voiture dans le réseau routier), l’analyse de grands jeux de données (réseaux sociaux, génomique, etc.). NVIDIA IndeX, une bibliothèque de visualisation de très grands jeux de données (répartis sur plusieurs GPU), est maintenant disponible comme extension à ParaView. cuDNN arrive en version 5, avec d’importantes améliorations de performance.

Sources : Pascal Architecture (images d’illustration), Nvidia Pascal Architecture & GP100 Specs Detailed – Faster CUDA Cores, Higher Thread Throughput, Smarter Scheduler & More, NVIDIA DGX-1 Pascal Based Supercomputer Announced – Features 8 Tesla P100 GPUs With 170 TFLOPs Compute, NVIDIA GTC 2016 Keynote Live Blog (transparents NVIDIA), Inside Pascal: NVIDIA’s Newest Computing Platform, CUDA 8 Features Revealed.

[SC15] AMD lance l’initiative Boltzmann

AMD s’est récemment lancé, comme Microsoft, dans une série d’actions d’ouverture de leur code source sous des licences libres, notamment au niveau de leurs pilotes pour Linux (AMDGPU). À un tout autre niveau, pour le calcul sur GPU, ils espèrent que leur architecture hétérogène (dite HSA) sera compatible avec les instructions de délégation de calcul d’OpenMP dans GCC 6 (qui devrait sortir début 2016), pour se mettre au même niveau qu’Intel (leur accélérateur Xeon Phi est déjà accessible par ce biais depuis GCC 5, c’est-à-dire début 2015). De même, ils explorent le côté LLVM des compilateurs libres, avec l’ouverture prévue du code de HCC, leur compilateur hétérogène pour leur plateforme HSA.

L’initiative Boltzmann prend le nom d’un physicien autrichien (ce qui n’est pas sans rappeler les noms de code des GPU NVIDIA), à l’origine de l’approche statistique en physique (ses travaux sont fondamentaux dans certaines utilisations actuelles des GPU). Cette initiative correspond à une revalorisation des GPU à destination des serveurs dans le marché du calcul de haute performance, avec notamment un pilote Linux prévu exclusivement pour le calcul sur ces GPU (sans aucune implémentation d’OpenGL). Leur compilateur HCC permettra d’y exécuter du code C ou C++ en utilisant OpenMP, un mécanisme de parallélisation assez général (pas initialement prévu pour les GPU), c’est-à-dire avec un seul et même langage et un seul compilateur pour une série de processeurs (de manière similaire au C++ AMP, proposé par Microsoft).

Une partie de cette initiative Boltzmann est prévue pour le portage des applications CUDA vers un « modèle de programmation C++ commun » aux différents types de processeurs disponibles, un modèle connu sous le doux nom de HIP (heterogeneous compute interface for portability). Il s’agit notamment d’effectuer une transpilation partielle du code CUDA, qui devrait être automatique pour nonante pour cent des cas courant — les dix pour cent restants devant être traduits à la main, ce décompte ne tenant pas compte de l’utilisation d’assembleur (PTX) ou de l’appel direct au pilote ; ce code transpilé sera toujours compilable pour les GPU NVIDIA. Au contraire, des applications CUDA compilées ne pourront pas directement être lancées sur un GPU AMD, ce qui nécessiterait l’implémentation complète dans le pilote d’une pile CUDA (et, accessoirement, une licence de la part de NVIDIA pour ce faire, déjà proposée dans le passé). Ce mouvement est absolument requis pour qu’AMD se relance dans la course du calcul scientifique de haute performance avec ses solutions GPGPU (et pas simplement des APU), au vu de la quantité de code CUDA existant.

Globalement, cette initiative Boltzmann matérialise un véritable retour en grande pompe dans le domaine du HPC, une niche très lucrative : le matériel existe déjà, mais l’environnement logiciel était encore défaillant pour reprendre des parts de marché, à Intel et NVIDIA. Les premiers résultats devraient arriver au premier trimestre 2016, avec des préversions. Restera à voir l’impact sur la performance.

Sources : AMD Launches ‘Boltzmann Initiative’ to Dramatically Reduce Barriers to GPU Computing on AMD FireProâ„¢ Graphics, AMD @ SC15: Boltzmann Initiative Announced – C++ and CUDA Compilers for AMD GPUs (images), AMD Plans To Contribute Heterogeneous Compute Compiler, AMD Working On CUDA Source Translation Support To Execute On FirePro GPUs.

CUDA et LLVM, des nouvelles

Le compilateur officiel de NVIDIA pour ses extensions GPGPU aux langages C et C++, nvcc, n’est plus la seule option pour compiler ce code CUDA depuis la version 4.1. Depuis ces quelques années, cette contribution au compilateur libre LLVM fait son bout de chemin. L’idée est de générer du code PTX, c’est-à-dire l’assembleur utilisé sur les GPU NVIDIA, traduit par les pilotes officiels en instructions pour chaque GPU.

La semaine dernière, deux nouveautés sont arrivées dans les dépôts de LLVM et Clang au sujet de CUDA. La première est un guide rapide pour compiler le compilateur et commencer à l’utiliser, surtout à destination de chercheurs et développeurs de LLVM, notamment pour les aider à développer des passes d’optimisation spécifiques aux GPU. Pour ce faire, une série de modifications au niveau des sources de LLVM est nécessaire (l’implémentation s’oriente comme nvcc, avec du code pour l’hôte et le GPU mélangé dans un même fichier).

Ce document explicite notamment les optimisations déjà réalisées par LLVM et les modifications CUDA pour améliorer la performance sur les GPU. Notamment :

  • la réduction des redondances dans le calcul d’expressions mathématiques : si x=b*s a déjà été exécuté, l’affectation y=(b+1)*s sera réécrite comme y=x+s, ce qui évite de calculer le facteur b+1 en sus de la multiplication ;
  • l’inférence de l’espace mémoire à utiliser pour une variable (à choisir entre mémoires globale, constante, partagée et locale) ;
  • le déroulement des boucles et l’extension en ligne des appels de fonction plus agressifs que sur du code à destination de CPU, puisque tout transfert de l’exécution sur un GPU est bien plus coûteux que sur un CPU (tout saut, toute condition nuit fortement à la performance) ;
  • l’analyse du recouvrement entre pointeurs, dans le cadre d’une passe généralisée à tout LLVM (bien que cette partie spécifique aux GPU n’est pas encore intégrée) ;
  • l’évitement de divisions lentes sur soixante-quatre bits, par manque de circuits spécifiques pour cette opération sur les GPU NVIDIA (elles sont remplacées, autant que possible, par des opérations sur trente-deux bits).

Une bonne partie de ces optimisations a eu des contributions de la part de Google, qui pourraient sembler quelque peu inattendues. Leur objectif est de développer une plateforme à la pointe de la technologie pour la recherche sur la compilation pour GPU et le calcul scientifique de haute performance en général. Leurs premiers résultats, présentés par Jingyue Wu, indiquent que leurs avancées sur LLVM font aussi bien que le compilateur de NVIDIA, nvcc, sur certains tests… et vont parfois jusque cinquante pour cent plus vite que le code produit par nvcc ! Ceci s’accompagne d’un gain de huit pour cent en moyenne sur le temps de compilation (sans véritable intégration à Clang, puisque la compilation s’effectue en plusieurs phases bien séparées — GPU et CPU —, ce qui fait perdre pas mal de temps).

Sortie de NVIDIA Nsight Visual Studio Edition 5.0

NVIDIA Nsight est une extension à Visual Studio (ou Eclipse) lui ajoutant diverses fonctionnalités dans le domaine du calcul sur GPU au sens large : l’outil est prévu pour des technologies de calcul pur (CUDA, OpenCL, DirectCompute) et de rendu 3D (Direct3D, OpenGL).

Cette nouvelle version, estampillée 5.0, est maintenant compatible avec Visual Studio 2015 (sauf pour les applications CUDA 7.5 !). Côté graphique, le débogage fonctionne également avec Direct3D 12, même si la majorité des fonctionnalités habituelles ne sont pas encore disponibles pour cette version (principalement celles axées sur l’analyse de la performance), mais devraient l’être dans une prochaine version. Les applications OpenG 4.4 et 4.5 sont aussi mieux gérées.

Pour le calcul proprement dit, les instantanés produits par le pilote en cas de plantage peuvent être ouverts dans Visual Studio pour une analyse plus détaillée. Les outils d’analyse sont désormais compatibles avec OpenCL 1.2. (NVIDIA étant largement en retard à ce niveau, AMD ayant déjà publié des pilotes compatibles OpenCL 2.0) Il devient également possible de déboguer et de profiler des applications CUDA distantes (à condition d’utiliser CUDA 7.5). Les systèmes d’exploitation 32 bits ne sont plus pris en charge et les processeurs de génération Fermi (GeForce 400/500, distribués jusqu’en 2012) ne le seront plus à la prochaine version.

Sources : Nsight Visual Studio Edition 5.0 New Features , Release Notes.

Sortie de CUDA 7.5 RC

La première préversion publique de CUDA 7.5 est disponible pour tous les développeurs CUDA enregistrés auprès de NVIDIA. La principale nouveauté est la gestion des nombres à virgule flottante codés sur seize bits, c’est-à-dire la moitié de l’encodage traditionnel (float sur trente-deux bits). Cette format est principalement utilisée pour du stockage de données quand la précision requise n’est pas très importante, mais aussi pour des calculs sous la même hypothèse.

Calculs sur seize bits

L’effet sur la performance du code peut être énorme : la bande passante requise est divisée par deux, ce qui permet de transmettre deux fois plus de nombres par unité de temps sur les bus existants, d’en stocker deux fois plus sur la même quantité de mémoire. Ils pourront donc se révéler très utiles pour les applications où ces éléments sont limitatifs, comme l’apprentissage de réseaux neuronaux de grande taille ou le filtrage de signaux en temps réel.

L’avantage en temps de calcul n’est disponible que sur les GPU ayant une partie prévue pour l’arithmétique sur seize bits, ce qui n’est pas le cas pour la majorité des processeurs disponibles, sauf sur Tegra. La prochaine architecture de GPU de NVIDIA, connue sous le nom de Pascal, aura des transistors alloués pour les calculs sur seize bits : le nombre de calculs effectués par seconde doublera entre les précisions FP16 et FP32 (le même rapport qu’entre FP32 et FP64). Entre temps, les calculs seront effectués en interne sur les mêmes circuits que précédemment, avec une précision bien plus élevée que des circuits dédiés (bien qu’elle soit en grande partie perdue lors de l’arrondi vers les seize bits).

Réseaux neuronaux

L’un des chevaux de bataille actuels de NVIDIA est l’apprentissage automatique, en particulier par réseaux neuronaux, plus spécifiquement profonds, par exemple pour une utilisation dans les voitures intelligentes. Ils ont par exemple développé la bibliothèque cuDNN (CUDA deep neural network), qui accélère les calculs par un GPU. La troisième version de cette bibliothèque a été optimisée, particulièrement au niveau des convolutions (FFT et 2D), ce qui améliore la performance lors de l’entraînement de réseaux neuronaux (gain d’un facteur deux pour l’apprentissage sur des GPU Maxwell). Elle gère également les nombres en virgule flottante sur seize bits, ce qui est utile pour des réseaux très grands, mais n’améliore pas (encore) les temps de calcul.

Améliorations de performance de cuDNN 3

NVIDIA a aussi développé l’outil DIGITS, un logiciel de bien plus haut niveau que cuDNN pour les mêmes réseaux neuronaux, prévu pour des profils plus scientifiques que pour des programmeurs. L’une des nouveautés est l’apprentissage distribué sur plusieurs GPU : ajouter un deuxième GPU aide à réduire fortement les temps de calcul (d’un facteur légèrement inférieur à deux), nettement moins impressionnant en ajoutant deux autres (facteur de deux et demi). Le gain sera probablement plus important avec les prochaines architectures de GPU, Pascal devant utiliser la technologie NVLink au lieu du bus PCI Express (partagé) pour la communication entre cartes.

DIGITS

Algèbre linéaire creuse

CUDA vient également avec la bibliothèque cuSPARSE pour l’algèbre linéaire sur des matrices creuses accélérée sur GPU. Une nouvelle opération vient d’y être ajoutée, nommée GEMVI, utilisée pour la multiplication entre une matrice pleine et un vecteur creux — la sortie étant évidemment un vecteur plein. Ce genre d’opérations est très utile pour l’apprentissage automatique, plus particulièrement dans le cas du traitement des langues. En effet, dans ce cas, un document rédigé dans une langue quelconque (français, anglais, allemand…) peut être représenté comme un comptage des occurrences de mots d’un dictionnaire ; bien évidemment, tous les mots du dictionnaire (même partiel) ne sont pas présents dans le texte, sa représentation vectorielle contient donc un grand nombre de zéros, il est donc creux. Une fois le dictionnaire défini, pour améliorer l’efficacité des traitements, le dictionnaire peut être réduit en taille pour n’en garder qu’un sous-espace vectoriel qui préserve la sémantique des textes : la transformation de la représentation du texte demande justement un produit entre le vecteur creux initial et une matrice de transformation.

C++11 et fonctions anonymes

La version précédente de CUDA a commencé à comprendre C++11, la dernière itération du langage de programmation. Les fonctions anonymes (lambdas) en font partie et servent notamment à écrire du code plus concis. CUDA 7.0 ne les tolérait que dans le code exécuté côté client, pas encore sur le GPU : ce point est corrigé, mais seulement comme fonctionnalité expérimentale. Par exemple, un code comptant les fréquences de quatre lettres dans une chaîne de caractères pourra s’écrire comme ceci :

void xyzw_frequency_thrust_device(int *count, char *text, int n) {
  using namespace thrust;

  *count = count_if(device, text, text+n, [] __device__ (char c) {
    for (const auto x : { 'x','y','z','w' })
      if (c == x) return true;
    return false;
  });
}

La même fonctionnalité permet d’écrire des boucles for pour une exécution en parallèle, avec une syntaxe similaire à OpenMP, par exemple une somme de deux vecteurs (SAXPY pour BLAS) :

void saxpy(float *x, float *y, float a, int N) {
    using namespace thrust;
    auto r = counting_iterator(0);
    for_each(device, r, r+N, [=] __device__ (int i) {
        y[i] = a * x[i] + y[i];
    });
}

Profilage

La dernière nouveauté annoncée concerne le profilage de code, nécessaire pour déterminer les endroits où les efforts d’amélioration de la performance doivent être investis en priorité. CUDA 7.5 améliore les outils NVIDIA Visual Profiler et NSight Eclipse Edition en proposant un profilage au niveau de l’instruction PTX (uniquement sur les GPU Maxwell GM200 et plus récents), pour détermines les lignes précises dans le code qui causent un ralentissement. Précédemment, le profilage ne pouvait se faire qu’au niveau d’un noyau, équivalent d’une fonction pour la programmation sur GPU (temps pris par le noyau à l’exécution, importance relative par rapport à l’exécution complète).

CUDA 6 avait déjà amélioré la situation en affichant une corrélation entre les lignes de code et le nombre d’instructions correspondantes. Cependant, un grand nombre d’instructions n’indique pas forcément que le noyau correspondant prendra beaucoup de temps à l’exécution. Pour remonter jusqu’à la source du problème, ces informations sont certes utiles, mais pas suffisantes, à moins d’avoir une grande expérience. Grâce à CUDA 7.5, le profilage se fait de manière beaucoup plus traditionnelle, avec un échantillonnage de l’exécution du programme, pour trouver les lignes qui prennent le plus de temps.

Sources et crédit des images : New Features in CUDA 7.5, NVIDIA @ ICML 2015: CUDA 7.5, cuDNN 3, & DIGITS 2 Announced.

Nouveau, le pilote NVIDIA libre pour Linux, s’ouvre à CUDA

Vulkan est prévue comme la prochaine évolution d’OpenGL, visant principalement à la rendre plus bas niveau, pour exploiter au mieux les possibilités des processeurs graphiques modernes (voir Vulkan : la nouvelle bibliothèque de hautes performances pour le GPU pour plus de détails). L’une de ses améliorations est SPIR-V, une représentation intermédiaire des shaders, une forme d’assembleur indépendant de la carte graphique utilisé pour représenter tout le code qu’un utilisateur peut y envoyer. Cette représentation intermédiaire sera partagée par OpenCL, prévu pour faciliter les calculs génériques sur GPU (comme la simulation de fluides), à l’instar de CUDA, une technologie propriétaire de NVIDIA, antérieure à OpenCL mais à l’évolution plus rapide.

Dans tout ce beau monde, peu d’annonces du côté du logiciel libre, si ce n’est un pilote pour les puces Intel développé par Valve et qui devrait être libéré. Récemment, Pierre Moreau a annoncé avoir commencé à travailler sur une passe de compilation SPIR-V vers la représentation intermédiaire des GPU NVIDIA NV50, déjà relativement anciens (série GeForce 8000, de 2006-2007). Cette étape est l’une des premières requises pour exécuter du code SPIR-V sur ces GPU.

En comptant les efforts actuels pour développer une passe de compilation LLVM IR vers SPIR-V, grâce à laquelle tout langage que LLVM et ses projets annexes (dont Clang) peuvent compiler pourrait être utilisé pour les GPU. En particulier, le code CUDA pourrait être transformé de la sorte et être exécuté sur Nouveau.

NVIDIA explore une autre piste pour exécuter du code CUDA grâce au pilote libre Nouveau. Leur intérêt pour ce pilote n’est pas neuf, mais il s’est focalisé sur leurs systèmes monopuces (SoC) Tegra, négligeant leurs autres GPU. Tout comme les efforts de Pierre Moreau, rien n’est actuellement visible au grand public en ce qui concerne le code ; cette fois, le développeur de NVIDIA est plus détaillé, en demandant des avis sur des modifications à apporter au code de la gestion de la mémoire du GPU.

Plus techniquement, il explique que CUDA a besoin d’avoir les mêmes adresses virtuelles sur le GPU et sur le CPU, ce que Nouveau ne permet pas encore. De plus, une application CUDA traditionnelle a tendance à allouer une grande quantité de mémoire en avance, pour une utilisation plus tardive ; des morceaux de cet espace sont alors alloués au fur et à mesure des besoins à l’exécution. Nouveau fusionne les opérations d’allocation des adresses virtuelles et de correspondance de mémoire, qui devraient être découplées pour ces utilisations.

Sources : Someone Is Already Working On SPIR-V For The Nouveau Driver, [Nouveau] Will nouveau support for cuda?, [LLVMdev] [RFC] Proposal for Adding SPIRV Target, NVIDIA Has Someone Working On Nouveau CUDA Support

Impact des GPU sur la simulation de fluides avec PyFR

La simulation de fluides (plus connue sous le sigle anglophone CFD, computational fluid dynamics) apparaît dans un grand nombre d’applications d’ingénierie : il s’agit de déterminer, par ordinateur, par exemple, l’écoulement de l’air autour d’avions ou de voitures de formule 1. Ainsi, les concepteurs peuvent adapter leurs pièces, notamment pour en améliorer l’aérodynamisme. Les calculs requis sont extrêmement poussés (équations aux dérivées partielles non linéaires) et requièrent une grande puissance de calcul.

Un grand nombre de logiciels existe déjà pour résoudre ces problèmes, comme OpenFOAM dans le monde du logiciel libre. Cependant, ils n’exploitent pas tous les accélérateurs disponibles pour ces calculs (OpenFOAM nécessite des extensions pour ce faire), comme les processeurs graphiques (GPU), à la mode dans le monde du calcul scientifique (HPSC) : l’intérêt est de bénéficier d’un très grand nombre de cÅ“urs. En effet, un processeur traditionnel a au plus une vingtaine de cÅ“urs par processeur : par exemple, dix-huit chez Intel). Un GPU en compte désormais plusieurs milliers (tant chez NVIDIA que AMD), bien qu’avec une architecture radicalement différente.

PyFR est un logiciel de CFD assez récent (sa version 1.0.0 est sortie la semaine dernière), basé sur un concept différent des solveurs actuels, la reconstruction de flux (FR), pour atteindre une précision bien plus élevée que les solveurs existants, même sur des géométries très complexes, tout en étant mieux adaptés à des GPU. (Techniquement, pour la discrétisation spatiale, il s’agit d’un schéma d’intégration à ordre élevé sur une grille non structurée, qui mêle la précision de méthodes d’ordre élevé des méthodes à différences finies et l’adaptabilité géométrique des volumes et éléments finis.)

Les solveurs actuels partent d’hypothèses des années 1980, quand les opérations de calcul sur des nombres réels étaient très coûteuses, mais la mémoire très rapide par rapport au processeur, deux points complètement dépassés par les architectures actuelles, afin de simuler des écoulements stationnaires. Ainsi, pour s’adapter à des simulations non stationnaires, il faut réinventer une série de composants

En pratique, ce nouveau logiciel peut donner des résultats dix fois plus précis dix fois plus rapidement que les techniques précédentes, tout en étant capable d’exploiter la puissance d’une série de GPU. Il s’adapte également à une série de plateformes : des grappes de CPU, formant les superordinateurs les plus courants ; des GPU, peu importe leur fabricant (tant AMD que NVIDIA). La distribution du calcul sur plusieurs nÅ“uds s’effectue par MPI.

Côté technologique, la parallélisation sur un nÅ“ud de calcul est possible tant par OpenMP (plusieurs fils d’exécution sur le même processeur) que CUDA ou OpenCL (GPU). Le logiciel lui-même est principalement codé en Python — ce qui ne l’empêche pas d’exceller au niveau des temps de calcul, grâce à la facilité d’appel de code natif, principalement en C et Fortran. L’interpréteur ajoute un surcoût en temps d’exécution inférieur au pour cent : cette perte de performance est négligeable quand le code de calcul ne compte que cinq mille lignes de code !

Tout le côté numérique est extrait dans une série de primitives bien comprises et optimisées. La communication s’effectue par une version de MPI adaptée à du calcul sur GPU, notamment en utilisant la technologie NVIDIA GPUDirect, qui permet la copie de données directement dans la mémoire du GPU, sans passer par la mémoire centrale (RAM). Les surcoûts dus à Python sont effacés en très grande partie grâce à la nature asynchrone de CUDA : le code Python n’est pas obligé d’attendre la fin des calculs sur le GPU avant de passer à autre chose.

Sources : On a Wing and PyFR: How GPU Technology Is Transforming Flow Simulation, PyFR: An open source framework for solving advection–diffusion type problems on streaming architectures using the flux reconstruction approach, PyFR: Technical Challenges of Bringing Next Generation Computational Fluid Dynamics to GPU Platforms – See more at: http://on-demand-gtc.gputechconf.com/gtcnew/on-demand-gtc.php?searchByKeyword=pyfr&searchItems=&sessionTopic=&sessionEvent=&sessionYear=&sessionFormat=&submit=&select=#sthash.CNa7viRf.dpuf (dont l’image).

Voir aussi : le site officiel de PyFR.

CUDA 7 RC

CUDA 7, dont la RC vient de sortir, est présentée par NVIDIA comme une « mise à jour importante de la plateforme ». Outre les habituelles améliorations de performances, du côté des plateformes, l’architecture IBM POWER sera maintenant gérée dans toutes les versions du SDK (seule la version 5.5 était disponible) — ce qui n’est pas étonnant, NVIDIA étant membre de la fondation OpenPOWER.

Lire la suite

Les prochains processeurs NVIDIA

Le CEO de NVIDIA, Jen-Hsun Huang, après avoir annoncé quelques détails architecturaux sur le futur des Tegra, a ébauché un futur où les technologies de la marque se rejoignent (mobile, desktop et superordinateurs). L’année dernière, l’architecture Kepler pour les GPU a été mise sur le marché, utilisée notamment pour des superordinateurs bien plus efficaces, énergétiquement parlant (voir le top 500 de novembre 2012).

La prochaine génération de GPU, Maxwell, offrira une mémoire virtuelle unifiée, donnant accès à la mémoire des GPU pour les CPU – et vice-versa. La suivante, Volta, tentera d’optimiser les accès à la mémoire sur le GPU, en plaçant les modules de mémoire juste au-dessus des puces de calcul (NVIDIA parle de mémoire empilée – « stacked DRAM ») ; les ingénieurs estiment une vitesse d’un téraoctet par seconde, soit cinquante disques Blu-Ray par seconde.

Planning prévisionnel pour les GPU NVIDIA.

Côté mobile, Logan mixera un processeur Tegra semblable aux actuels avec un GPU Kepler, sur la même puce, dans les mêmes dimensions qu’un Tegra actuel. Parker ajoutera le support d’ARM 64 bits, en plus de passer à un GPU Maxwell.

Planning prévisionnel pour les processeurs NVIDIA Tegra.

Source : http://blogs.nvidia.com/2013/03/nvidia-ceo-updates-nvidias-roadmap/