Tout le monde sait depuis longtemps que sur les cartes vidéo, vous pouvez non seulement jouer à des jouets, mais aussi effectuer des choses qui ne sont pas liées aux jeux, par exemple, former un réseau de neurones, vous souvenir de la crypto-monnaie ou effectuer des calculs scientifiques. Comment cela s'est passé, vous pouvez le lire ici , mais je voulais aborder la question de savoir pourquoi le GPU peut être intéressant pour le programmeur moyen (non lié à GameDev) comment aborder le développement sur le GPU sans y consacrer beaucoup de temps, décider si regardez dans cette direction, et " déterminez sur vos doigts" quel profit vous pouvez obtenir.

L'article a été écrit sur la base de ma présentation en HighLoad ++. Il aborde principalement les technologies proposées par NVIDIA. Je n'ai aucun but de faire la publicité de produits, je les donne simplement à titre d'exemple, et à coup sûr, quelque chose de similaire peut être trouvé chez les fabricants concurrents.

Pourquoi compter sur le GPU?

Deux processeurs peuvent être comparés selon différents critères, les plus populaires étant probablement la fréquence et le nombre de cœurs, la taille des caches, etc., mais au final, nous nous intéressons au nombre d'opérations qu'un processeur peut effectuer par unité de temps, de quel type d'opération il s'agit, mais une question distincte Une métrique commune est le nombre d'opérations en virgule flottante par seconde - flops. Et lorsque nous voulons comparer le chaud au doux, et dans notre cas le GPU avec le CPU, cette métrique est très pratique.

Le graphique ci-dessous montre la croissance de ces mêmes flops au fil du temps pour les processeurs et les cartes vidéo.

(Les données sont collectées à partir de sources ouvertes, il n'y a pas de données pour 2019-20 ans, car tout n'est pas si beau là-bas, mais les GPU gagnent toujours)

Eh bien, c'est tentant, n'est-ce pas? Nous déplaçons tous les calculs du CPU vers le GPU et obtenons huit fois les meilleures performances!

Mais, bien sûr, tout n'est pas si simple. Vous ne pouvez pas tout prendre et tout transférer sur le GPU, pourquoi, nous parlerons plus loin.

Architecture GPU et sa comparaison avec le CPU

J'apporte à beaucoup une image familière de l'architecture du CPU et des éléments de base:

CPU Core

Qu'est-ce qui est si spécial? Un cœur et un tas de blocs auxiliaires.

Voyons maintenant l'architecture GPU:

GPU Core

Une carte vidéo a beaucoup de cœurs de traitement, généralement plusieurs milliers, mais ils sont combinés en blocs; pour les cartes vidéo NVIDIA, généralement 32 chacune, et ont des éléments communs, y compris et registres. L'architecture du noyau GPU et des éléments logiques est beaucoup plus simple que sur le CPU, à savoir, il n'y a pas de prefetchers, de prédicteurs de brunch et bien plus encore.

Eh bien, ce sont les points clés de la différence dans l'architecture du CPU et du GPU, et, en fait, ils imposent des restrictions ou, inversement, ouvrent les possibilités de ce que nous pouvons lire efficacement sur le GPU.

Je n'ai pas mentionné un autre point important, généralement la carte vidéo et le processeur ne «fouinent» pas entre eux et n'écrivent pas de données sur la carte vidéo et ne lisent pas le résultat - ce sont des opérations distinctes et peuvent se révéler être un «goulot d'étranglement» dans votre système, un graphique du temps de pompage par rapport à la taille les données sont données plus loin dans l'article.

Limitations et fonctionnalités du GPU

Quelles limitations cette architecture impose-t-elle aux algorithmes exécutables:

Si nous calculons sur un GPU, alors nous ne pouvons pas sélectionner un seul cœur, un bloc entier de cœurs sera alloué (32 pour NVIDIA).
Tous les cœurs exécutent les mêmes instructions, mais avec des données différentes (nous en parlerons plus tard), ces calculs sont appelés Single-Instruction-Multiple-Data ou SIMD (bien que NVIDIA introduise son raffinement).
En raison de l'ensemble relativement simple de blocs logiques et de registres généraux, le GPU n'aime vraiment pas la ramification, et en fait la logique complexe dans les algorithmes.

Quelles opportunités ouvre-t-il:

En fait, l'accélération de ces mêmes calculs SIMD. L'exemple le plus simple est l'ajout élémentaire de matrices, et analysons-le.

Réduction des algorithmes classiques à la représentation SIMD

Transformation

Nous avons deux tableaux, A et B, et nous voulons ajouter un élément du tableau B à chaque élément du tableau A. Ci-dessous est un exemple en C, bien que j'espère qu'il sera clair pour ceux qui ne parlent pas ce langage:

void func(float *A, float *B, size)
{ 
   for (int i = 0; i < size; i++) 
   { 
       A[i] += B[i]
   } 
}

Bouclage classique des éléments dans une boucle et exécution linéaire.

Voyons maintenant à quoi ressemblera ce code pour le GPU:

void func(float *A, float *B, size) 
{ 
   int i = threadIdx.x; 
   if (i < size) 
      A[i] += B[i] 
}

Et ici, c'est déjà intéressant, la variable threadIdx est apparue, que nous ne semblions déclarer nulle part. Oui, son système nous fournit. Imaginez que dans l'exemple précédent, le tableau se compose de trois éléments et que vous souhaitez l'exécuter dans trois threads parallèles. Pour ce faire, vous devez ajouter un autre paramètre - le numéro d'index ou de flux. C'est ce que la carte vidéo fait pour nous, bien qu'elle passe l'index en tant que variable statique et peut fonctionner avec plusieurs dimensions à la fois - x, y, z.

Autre nuance, si vous souhaitez démarrer un grand nombre de flux parallèles à la fois, les flux devront être divisés en blocs (une caractéristique architecturale des cartes vidéo). La taille de bloc maximale dépend de la carte vidéo, et l'indice de l'élément pour lequel nous effectuons les calculs devra être obtenu comme suit:

int i = blockIdx.x * blockDim.x + threadIdx.x; // blockIdx –  , blockDim –  , threadIdx –

En conséquence, ce que nous avons: beaucoup de threads exécutés en parallèle qui exécutent le même code, mais avec des indices différents, et, par conséquent, des données, c'est-à-dire le même SIMD.

C'est l'exemple le plus simple, mais si vous souhaitez travailler avec le GPU, vous devez amener votre tâche sous la même forme. Malheureusement, ce n'est pas toujours possible et dans certains cas peut faire l'objet d'une thèse de doctorat, mais néanmoins, les algorithmes classiques peuvent encore être amenés à cette forme.

Agrégation

Voyons maintenant à quoi ressemblera l'agrégation castée en représentation SIMD:

Nous avons un tableau de n éléments. À la première étape, nous démarrons n / 2 threads et chaque thread ajoute deux éléments, à savoir en une seule itération, nous additionnons la moitié des éléments du tableau. Et puis dans la boucle, nous répétons la même chose pour le tableau nouvellement créé, jusqu'à ce que nous agrégions les deux derniers éléments. Comme vous pouvez le voir, plus la taille du tableau est petite, moins nous pouvons démarrer de threads parallèles, c'est-à-dire sur un GPU, il est logique d'agréger des tableaux d'une taille suffisamment grande. Un tel algorithme peut être utilisé pour calculer la somme des éléments (en passant, n'oubliez pas le débordement possible du type de données avec lequel vous travaillez), rechercher le maximum, le minimum ou simplement la recherche.

Tri

Mais le tri semble déjà beaucoup plus compliqué.

Les deux algorithmes de tri les plus populaires sur le GPU sont:

Tri bitonique
Radix-sort

Mais radix-sort est toujours utilisé plus souvent, et une implémentation prête pour la production peut être trouvée dans certaines bibliothèques. Je n'analyserai pas en détail le fonctionnement de ces algorithmes; ceux qui sont intéressés peuvent trouver une description de radix-sort sur https://www.codeproject.com/Articles/543451/Parallel-Radix-Sort-on-the-GPU-using-Cplusplus- AMP et https://stackoverflow.com/a/26229897

Mais l'idée est que même un algorithme non linéaire tel que le tri peut être réduit à une vue SIMD.

Et maintenant, avant de regarder les vrais chiffres qui peuvent être obtenus à partir du GPU, essayons de comprendre comment programmer pour ce miracle de la technologie?

Où commencer

Les deux technologies les plus courantes pouvant être utilisées pour la programmation sous le GPU:

Opencl
Cuda

OpenCL est une norme prise en charge par la plupart des fabricants de cartes vidéo, notamment et sur les appareils mobiles, le code écrit en OpenCL peut également être exécuté sur le CPU.

Vous pouvez utiliser OpenCL à partir de C / C ++, il existe des liants vers d'autres langages.

Pour OpenCL, j'ai le plus apprécié le livre OpenCL en action . Il décrit également différents algorithmes sur le GPU, notamment Tri bitonique et tri Radix.

CUDA est la technologie et le SDK exclusifs de NVIDIA. Vous pouvez écrire en C / C ++ ou utiliser des liaisons vers d'autres langages.

Comparer OpenCL et CUDA n'est pas correct, car l'un est la norme, l'autre est l'ensemble du SDK. Néanmoins, beaucoup de gens choisissent CUDA pour le développement de cartes vidéo, malgré le fait que la technologie est propriétaire, bien que gratuite et ne fonctionne que sur les cartes NVIDIA. Il y a plusieurs raisons à cela:

API
, GPU, (host)
, ..

Les particularités incluent le fait que CUDA est livré avec son propre compilateur, qui peut également compiler du code C / C ++ standard.

Le livre CUDA le plus complet que j'ai rencontré était Professional CUDA C Programming , bien qu'il soit déjà un peu dépassé, il aborde néanmoins de nombreuses nuances techniques de programmation pour les cartes NVIDIA.

Mais que faire si je ne veux pas passer quelques mois à lire ces livres, écrire mon propre programme pour une carte vidéo, tester et déboguer, puis découvrir que ce n'est pas pour moi?

Comme je l'ai dit, il existe un grand nombre de bibliothèques qui cachent la complexité du développement sous le GPU: XGBoost, cuBLAS, TensorFlow, PyTorch et autres, nous considérerons la bibliothèque de poussée, car il est moins spécialisé que les autres bibliothèques ci-dessus, mais en même temps il implémente des algorithmes de base, par exemple, le tri, la recherche, l'agrégation et, avec une forte probabilité, il peut être applicable à vos tâches.

Thrust est une bibliothèque C ++ qui vise à "remplacer" les algorithmes STL standard par des algorithmes basés sur GPU. Par exemple, le tri d'un tableau de nombres à l'aide de cette bibliothèque sur une carte vidéo ressemblerait à ceci:

thrust::host_vector<DataType> h_vec(size); //    
std::generate(h_vec.begin(), h_vec.end(), rand); //   
thrust::device_vector<DataType> d_vec = h_vec; //         
thrust::sort(d_vec.begin(), d_vec.end()); //    
thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin()); //   ,

(n'oubliez pas que l'exemple doit être compilé par un compilateur de NVIDIA)

Comme vous pouvez le voir, thrust :: sort est très similaire à un algorithme similaire de STL. Cette bibliothèque cache de nombreuses difficultés, en particulier le développement d'un sous-programme (plus précisément, le noyau), qui sera exécuté sur la carte vidéo, mais en même temps prive de flexibilité. Par exemple, si nous voulons trier plusieurs gigaoctets de données, il serait logique d'envoyer une donnée à la carte pour commencer le tri, et pendant le tri, envoyez plus de données à la carte. Cette approche est appelée masquage de latence et permet une utilisation plus efficace des ressources de mappage de serveur, mais, malheureusement, lorsque nous utilisons des bibliothèques de haut niveau, ces opportunités restent cachées. Mais pour le prototypage et la mesure des performances, ce sont les mêmes, en particulier avec la poussée, vous pouvez mesurer les frais généraux fournis par le transfert de données.

J'ai écrit une petite référence en utilisant cette bibliothèque, qui exécute plusieurs algorithmes populaires avec différentes quantités de données sur le GPU, voyons quels sont les résultats.

Résultats de l'algorithme GPU

Pour tester le GPU, j'ai pris une instance dans AWS avec une carte vidéo Tesla k80, ce n'est pas la carte serveur la plus puissante à ce jour (la Tesla v100 la plus puissante), mais la plus abordable et embarquée:

4992 CUDA Kernels
24 Go de mémoire
480 Gb / s - bande passante mémoire

Et pour les tests sur le CPU, j'ai pris une instance avec un processeur Intel Xeon CPU E5-2686 v4 @ 2.30GHz

Transformation

Temps d'exécution de la transformation sur le GPU et le CPU en ms

Comme vous pouvez le voir, la transformation habituelle des éléments du tableau est approximativement la même dans le temps, à la fois sur le GPU et sur le CPU. Et pourquoi? Parce que le surcoût pour l'envoi de données vers la carte et le dos absorbe toute l'augmentation des performances (nous parlerons du surcoût séparément), et il y a relativement peu de calculs sur la carte. De plus, n'oubliez pas que les processeurs prennent également en charge les instructions SIMD et que les compilateurs dans des cas simples peuvent les utiliser efficacement.

Voyons maintenant l'efficacité de l'agrégation sur le GPU.

Agrégation

Temps d'exécution d'agrégation sur GPU et CPU en ms

Dans l'exemple d'agrégation, nous constatons déjà une augmentation significative des performances avec une augmentation du volume de données. Il convient également de prêter attention au fait que nous pompons une grande quantité de données dans la mémoire de la carte, et qu'une seule valeur agrégée est reprise, c'est-à-dire Les frais généraux pour le transfert de données de la carte vers la RAM sont minimes.

Passons à l'exemple le plus intéressant - le tri.

Tri

Temps de tri vers le GPU et le CPU en ms

Malgré le fait que nous envoyons l'intégralité du tableau de données à la carte vidéo et vice versa, le tri vers le GPU 800 Mo de données est environ 25 fois plus rapide que sur le processeur.

Frais généraux de transfert de données

Comme le montre l'exemple de transformation, il n'est pas toujours évident que le GPU sera efficace même dans les tâches qui sont bien parallèles. La raison en est un surcoût pour le transfert de données de la RAM de l'ordinateur vers la mémoire de la carte vidéo (dans les consoles de jeu, en passant, la mémoire est partagée entre le CPU et le GPU, et il n'est pas nécessaire de transférer des données). Une des caractéristiques d'une carte vidéo est la bande passante mémoire ou bande passante mémoire, qui détermine la bande passante théorique de la carte. Pour Tesla k80, elle est de 480 Go / s, pour Tesla v100, elle est déjà de 900 Go / s. De plus, la version PCI Express et l'implémentation de la façon dont vous transférerez les données sur la carte affecteront le débit, par exemple, cela peut être fait dans plusieurs flux parallèles.

Examinons les résultats pratiques obtenus pour la carte graphique Tesla k80 dans le cloud Amazon:

Temps de transfert des données vers le GPU, tri et transfert des données vers la RAM en ms

HtoD - transfert des données vers la carte vidéo

GPU Exécution - tri sur la carte vidéo

DtoH - copie des données de la carte vidéo vers la RAM

La première chose à noter est que la lecture des données de la carte vidéo est plus rapide que écrivez-les là-bas.

La seconde - lorsque vous travaillez avec une carte vidéo, vous pouvez obtenir une latence de 350 microsecondes, et cela peut déjà être suffisant pour certaines applications à faible latence.

Le graphique ci-dessous montre une surcharge pour plus de données:

Temps de transfert des données vers le GPU, tri et transfert des données vers la RAM en ms

Utilisation du serveur

La question la plus courante est de savoir en quoi une carte vidéo de jeu diffère d'une carte serveur? Selon les caractéristiques, ils sont très similaires, mais les prix diffèrent considérablement.

Les principales différences entre le serveur (NVIDIA) et la carte de jeu:

Garantie du fabricant (la carte de jeu n'est pas conçue pour une utilisation sur serveur)
Problèmes de virtualisation possibles pour une carte graphique grand public
Disponibilité du mécanisme de correction d'erreur sur la carte serveur
Le nombre de threads parallèles (pas les cœurs CUDA) ou la prise en charge d'Hyper-Q, qui vous permet de travailler avec la carte à partir de plusieurs threads sur le CPU, par exemple, télécharger des données sur la carte à partir d'un thread et démarrer les calculs à partir d'un autre

Ce sont peut-être les principales différences importantes que j'ai trouvées.

Multithreading

Après avoir compris comment exécuter l'algorithme le plus simple sur la carte vidéo et quels résultats peuvent être attendus, la prochaine question logique est de savoir comment la carte vidéo se comportera lors du traitement de plusieurs demandes parallèles. En réponse, j'ai deux graphiques de calcul sur le GPU et un processeur à 4 et 32 cœurs:

Temps nécessaire pour effectuer des calculs mathématiques sur le GPU et le CPU avec des matrices de 1000 x 60 en ms

. Ce graphique effectue des calculs avec des matrices de 1000 x 60 éléments. Les calculs sont lancés à partir de plusieurs flux de programme, un flux séparé est créé pour le GPU pour chaque flux CPU (le très Hyper-Q est utilisé).

Comme vous pouvez le voir, le processeur gère très bien cette charge, tandis que la latence pour une demande par GPU augmente considérablement avec une augmentation du nombre de demandes parallèles.

Le temps pour effectuer des calculs mathématiques sur le GPU et le CPU avec des matrices 10 000 x 60 en ms.

Sur le deuxième graphique, les mêmes calculs, mais avec des matrices 10 fois plus longues, et le GPU se comporte beaucoup mieux sous une telle charge. Ces graphiques sont très indicatifs, et nous pouvons conclure: le comportement sous charge dépend de la nature de la charge elle-même. Un processeur peut également gérer les calculs matriciels assez efficacement, mais dans une certaine mesure. Pour une carte vidéo, il est caractéristique que pour une petite charge de calcul, les performances chutent de façon approximativement linéaire. Avec une augmentation de la charge et du nombre de threads parallèles, la carte vidéo s'adapte mieux.

Il est difficile de supposer comment le GPU se comportera dans diverses situations, mais comme vous pouvez le voir, dans certaines conditions, une carte serveur peut traiter les demandes de plusieurs flux parallèles assez efficacement.

Nous discuterons de quelques autres questions que vous pourriez avoir si vous décidez toujours d'utiliser le GPU dans vos projets.

Limite de ressources

Comme nous l'avons déjà dit, les deux principales ressources d'une carte vidéo sont le calcul des cœurs et de la mémoire.

Par exemple, nous avons plusieurs processus ou conteneurs utilisant une carte vidéo, et nous aimerions pouvoir partager la carte vidéo entre eux. Malheureusement, il n'y a pas d'API simple pour cela. NVIDIA propose la technologie vGPU , mais je n'ai pas trouvé la carte Tesla k80 dans la liste des cartes prises en charge, et d'après ce que je peux comprendre de la description, la technologie est plus axée sur les écrans virtuels que sur les calculs. AMD propose peut-être quelque chose de plus approprié.

Par conséquent, si vous prévoyez d'utiliser le GPU dans vos projets, vous devez vous fier au fait que l'application utilisera exclusivement la carte vidéo, ou vous contrôlerez par programme la quantité de mémoire allouée et le nombre de cœurs utilisés pour les calculs.

Conteneurs et GPU

Si vous avez déterminé la limite de ressources, alors la question logique suivante: que faire s'il y a plusieurs cartes vidéo sur le serveur?

Encore une fois, vous pouvez décider au niveau de l'application quel GPU il utilisera.

Les conteneurs Docker sont un autre moyen plus pratique. Vous pouvez utiliser des conteneurs réguliers, mais NVIDIA propose ses conteneurs NGC , avec des versions optimisées de divers logiciels, bibliothèques et pilotes. Pour un conteneur, vous pouvez limiter le nombre de GPU utilisés et leur visibilité sur le conteneur. Les frais généraux liés à l'utilisation des conteneurs sont d'environ 3%.

Travailler en cluster

Une autre question, que faire si vous souhaitez effectuer une tâche sur plusieurs GPU au sein du même serveur ou cluster?

Si vous avez choisi une bibliothèque similaire à Thrust ou une solution de niveau inférieur, la tâche devra être résolue manuellement. Les cadres de haut niveau, tels que pour l'apprentissage automatique ou les réseaux de neurones, prennent généralement en charge la possibilité d'utiliser plusieurs cartes prêtes à l'emploi.

De plus, je voudrais noter que, par exemple, NVIDIA propose une interface pour l'échange direct de données entre les cartes - NVLINK , qui est nettement plus rapide que PCI Express. Et il existe une technologie pour l'accès direct à la mémoire de la carte à partir d'autres périphériques PCI Express - GPUDirect RDMA , incl. et réseau .

Recommandations

Si vous envisagez d'utiliser le GPU dans vos projets, le GPU vous convient probablement si:

Votre tâche peut être réduite à une vue SIMD
Il est possible de charger la plupart des données sur la carte avant les calculs (cache)
Le défi passe par l'informatique intensive

Vous devez également poser des questions à l'avance:

Combien de requêtes parallèles seront
Quelle latence attendez-vous
Avez-vous besoin d'une carte pour votre charge? Avez-vous besoin d'un serveur avec plusieurs cartes ou d'un cluster de serveurs GPU

C'est tout, j'espère que le matériel vous sera utile et vous aidera à prendre la bonne décision!

Références

Benchmark et résultats sur github - https://github.com/tishden/gpu_benchmark/tree/master/cuda

En plus du sujet, un enregistrement du rapport «GPU Databases - Architecture, Performance and Prospects for Use»

NVIDIA NGC Containers Webinar Webinaires - http : //bit.ly/2UmVIVt ou http://bit.ly/2x4vJKF

GPU Computing - Pourquoi, quand et comment. Plus quelques tests

Pourquoi compter sur le GPU?

Architecture GPU et sa comparaison avec le CPU

Limitations et fonctionnalités du GPU

Réduction des algorithmes classiques à la représentation SIMD

Transformation

Agrégation

Tri

Où commencer

Résultats de l'algorithme GPU

Transformation

Agrégation

Tri

Frais généraux de transfert de données

Utilisation du serveur

Multithreading

Limite de ressources

Conteneurs et GPU

Travailler en cluster

Recommandations

Références

More articles: