Quels types de problèmes se prêtent bien à l’informatique GPU?

84

J'ai donc une tête décente pour les problèmes avec lesquels je travaille qui sont les meilleurs en série et qui peuvent être gérés en parallèle. Mais pour le moment, je n'ai pas une très bonne idée de ce qui est le mieux géré par un calcul basé sur le processeur, et de ce qui devrait être transféré à un GPU.

Je sais que c’est une question fondamentale, mais une bonne partie de mes recherches se trouve coincée entre des personnes qui militent clairement pour l’un ou l’autre sans vraiment justifier pourquoi , ou des règles empiriques quelque peu vagues. Vous cherchez une réponse plus utile ici.

Fomite
la source

Réponses:

63

Le matériel GPU présente deux points forts: le calcul brut (FLOP) et la bande passante mémoire. Les problèmes de calcul les plus difficiles tombent dans l'une de ces deux catégories. Par exemple, l'algèbre linéaire dense (A * B = C ou Résoudre [Ax = y] ou Diagonalize [A], etc.) se situe quelque part dans le spectre de la bande passante calcul / mémoire en fonction de la taille du système. Les transformées de Fourier rapides (FFT) s’adaptent également à ce moule avec des besoins élevés en bande passante globale. Comme d'autres transformations, algorithmes grille / maillage, Monte Carlo, etc. Si vous examinez les exemples de code NVIDIA SDK , vous pouvez avoir une idée du type de problème le plus souvent traité.

Je pense que la réponse la plus instructive est à la question «À quel genre de problèmes les GPU sont-ils vraiment mauvais? La plupart des problèmes qui ne font pas partie de cette catégorie peuvent être exécutés sur le processeur graphique, bien que certains nécessitent plus d'efforts que d'autres.

Les problèmes qui ne correspondent pas bien sont généralement trop petits ou trop imprévisibles. De très petits problèmes manquent du parallélisme nécessaire pour utiliser tous les threads du GPU et / ou pourraient entrer dans un cache de bas niveau sur le CPU, améliorant considérablement les performances du CPU. Les problèmes imprévisibles ont trop de branches significatives, ce qui peut empêcher les données de s'écouler efficacement de la mémoire du processeur graphique vers les cœurs ou réduire le parallélisme en brisant le paradigme SIMD (voir « Divergences distordues »). Voici des exemples de ce genre de problèmes:

  • La plupart des algorithmes de graphes (trop imprévisibles, en particulier dans l'espace mémoire)
  • Algèbre linéaire clairsemée (mais c'est aussi mauvais pour le processeur)
  • Petits problèmes de traitement du signal (FFT inférieure à 1000 points, par exemple)
  • Chercher
  • Trier
Max Hutchinson
la source
3
Néanmoins, des solutions GPU pour ces problèmes "imprévisibles" sont possibles et, bien que ce ne soit généralement pas faisable de nos jours, pourraient gagner en importance à l'avenir.
gauche du
6
Je voudrais ajouter spécifiquement des branches à la liste des disjoncteurs de performances GPU. Vous voulez que tous vos (des centaines) exécutent la même instruction (comme dans SIMD) pour effectuer un calcul réellement parallèle. Par exemple, sur les cartes AMD si l'un des flux d'instructions rencontre une branche et doit diverger - tout le front d'onde (groupe parallèle) diverge. Si les autres unités du front d'onde ne doivent pas diverger, elles doivent effectuer une seconde passe. C'est ce que maxhutch entend par prévisibilité, je suppose.
Violet Giraffe
2
@VioletGiraffe, ce n'est pas nécessairement vrai. Dans CUDA (c'est-à-dire sur les GPU Nvidia), la divergence de branche affecte uniquement la chaîne en cours, qui est au plus de 32 threads. Différentes chaînes, bien qu'elles exécutent le même code, ne sont pas synchrones à moins d'être explicitement synchronisées (par exemple avec __synchtreads()).
Pedro
1
@ Pedro: C'est vrai, mais la création de branches nuit en général aux performances. Pour les codes hautes performances (quel code GPU n'est pas?), Il est presque essentiel de prendre cela en compte.
jvriesem
21

Les problèmes ayant une intensité arithmétique élevée et des schémas d'accès à la mémoire réguliers sont généralement faciles à implémenter sur des GPU et donnent de bons résultats.

La difficulté fondamentale du code GPU hautes performances est que vous avez une tonne de cœurs et que vous voulez qu'ils soient tous utilisés au maximum de leur potentiel. Cela est difficile en raison de problèmes d’accès mémoire irrégulier ou d’intensité arithmétique élevée: vous passez beaucoup de temps à communiquer les résultats ou vous passez trop de temps à extraire des éléments de la mémoire (ce qui est lent!) Et vous ne prenez pas assez de temps pour calculer les chiffres. Bien entendu, le potentiel de simultanéité de votre code est essentiel pour sa bonne mise en œuvre sur le GPU.

Reid.Atcheson
la source
Pouvez-vous préciser ce que vous entendez par modèles d'accès mémoire classiques?
Fomite
1
La réponse de Maxhutch est meilleure que la mienne. Ce que je veux dire par schéma d’accès régulier, c’est que la mémoire est accédée de manière temporelle et localement. C'est-à-dire que vous ne faites pas d'énormes sauts autour de la mémoire à plusieurs reprises. C'est aussi un peu un forfait que j'ai remarqué. Cela signifie également que vos schémas d'accès aux données peuvent être prédéterminés soit par le compilateur, soit par vous, le programmeur, afin de minimiser les branchements (instructions conditionnelles dans le code).
Reid.Atcheson
15

Ceci n'est pas une réponse en soi, mais plutôt un ajout aux autres réponses de maxhutch et Reid.Atcheson .

Pour tirer le meilleur parti des GPU, votre problème doit non seulement être hautement parallèle (ou massivement), mais également l'algorithme principal à exécuter sur le GPU doit être aussi petit que possible. En termes OpenCL , il s’agit principalement du noyau .

Pour être plus précis, le noyau doit s’inscrire dans le registre de chaque unité de multitraitement (ou unité de calcul ) du GPU. La taille exacte du registre dépend du GPU.

Étant donné que le noyau est suffisamment petit, les données brutes du problème doivent tenir dans la mémoire locale du GPU (lecture: mémoire locale (OpenCL) ou mémoire partagée (CUDA) d'une unité de calcul). Sinon, même la bande passante mémoire élevée du GPU n’est pas assez rapide pour que les éléments de traitement soient occupés en permanence.
Habituellement , cette mémoire est d' environ 16 à 32 Kiocte grand .

Torbjörn
la source
La mémoire locale / partagée de chaque unité de traitement n'est-elle pas partagée entre toutes les dizaines (?) De threads s'exécutant dans un seul cluster de cœurs? Dans ce cas, n’avez-vous pas besoin de réduire considérablement la taille de votre ensemble de données afin de tirer le meilleur parti des performances du GPU?
Dan Neely
La mémoire locale / partagée d'une unité de traitement n'est accessible que par l'unité de calcul elle-même et n'est donc partagée que par les éléments de traitement de cette unité de calcul. La mémoire globale de la carte graphique (généralement 1 Go) est accessible à toutes les unités de traitement. La bande passante entre les éléments de traitement et la mémoire locale / partagée est très rapide (> 1 To / s) mais la bande passante vers la mémoire globale est beaucoup plus lente (~ 100 Go / s) et doit être partagée entre toutes les unités de calcul.
Torbjörn
Je ne parlais pas de la mémoire principale du GPU. Je pensais que la mémoire interne n'était allouée qu'au groupe de niveaux, pas par noyau. ex pour un nVidia GF100 / 110 gpu; pour chacun des 16 groupes SM, et non les 512 cuda cœurs. Avec chaque SM conçu pour exécuter jusqu'à 32 threads en parallèle, l'optimisation des performances du processeur graphique nécessiterait de maintenir le jeu de travail dans la plage de 1 Ko / thread.
Dan Neely
@Torbjoern Ce que vous voulez, c’est de garder tous les pipelines d’exécution de GPU occupés. Les GPU y parviennent de deux manières: (1) le moyen le plus courant consiste à augmenter l’occupation, ou autrement, en augmentant le nombre de threads simultanés les ressources partagées afin que vous puissiez avoir plus de threads actifs); peut-être mieux, consiste à (2) augmenter le parallélisme du niveau d’instruction au sein de votre noyau, de sorte que vous puissiez avoir un noyau plus gros avec une occupation relativement faible (petit nombre de threads actifs). See bit.ly/Q3KdI0
fcruz
11

Probablement un ajout plus technique aux réponses précédentes: les GPU CUDA (c.-à-d. Nvidia) peuvent être décrits comme un ensemble de processeurs fonctionnant de manière autonome sur 32 threads chacun. Les threads de chaque processeur fonctionnent en mode verrouillé (pensez à SIMD avec des vecteurs de longueur 32).

Bien que la façon la plus tentante de travailler avec des GPU soit de prétendre que tout fonctionne parfaitement, ce n’est pas toujours la manière la plus efficace de procéder.

Si votre code ne pas paralléliser bien / automatiquement des centaines / milliers de fils, vous pourriez être en mesure de le décomposer en tâches asynchrones individuelles qui ne bien paralléliser, et Execute ceux avec seulement 32 threads en cours d' exécution dans la serrure à pas. CUDA fournit un ensemble d'instructions atomiques qui permettent d'implémenter des mutex, ce qui permet aux processeurs de se synchroniser entre eux et de traiter une liste de tâches dans un paradigme de pool de threads . Votre code fonctionnerait alors sensiblement de la même manière que sur un système multicœur, gardez simplement à l'esprit que chaque cœur a alors 32 threads.

Voici un petit exemple, en utilisant CUDA, de la façon dont cela fonctionne

/* Global index of the next available task, assume this has been set to
   zero before spawning the kernel. */
__device__ int next_task;

/* We will use this value as our mutex variable. Assume it has been set to
   zero before spawning the kernel. */
__device__ int tasks_mutex;

/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
    while ( atomicCAS( m , 0 , 1 ) != 0 );
    }
__device__ inline void cuda_mutex_unlock ( int *m ) {
    atomicExch( m , 0 );
    }

__device__ void task_do ( struct task *t ) {

    /* Do whatever needs to be done for the task t using the 32 threads of
       a single warp. */
    }

__global__ void main ( struct task *tasks , int nr_tasks ) {

    __shared__ task_id;

    /* Main task loop... */
    while ( next_task < nr_tasks ) {

        /* The first thread in this block is responsible for picking-up a task. */
        if ( threadIdx.x == 0 ) {

            /* Get a hold of the task mutex. */
            cuda_mutex_lock( &tasks_mutex );

            /* Store the next task in the shared task_id variable so that all
               threads in this warp can see it. */
            task_id = next_task;

            /* Increase the task counter. */
            next_tast += 1;

            /* Make sure those last two writes to local and global memory can
               be seen by everybody. */
            __threadfence();

            /* Unlock the task mutex. */
            cuda_mutex_unlock( &tasks_mutex );

            }

        /* As of here, all threads in this warp are back in sync, so if we
           got a valid task, perform it. */
        if ( task_id < nr_tasks )
            task_do( &tasks[ task_id ] );

        } /* main loop. */

    }

Vous devez ensuite appeler le noyau avec main<<<N,32>>>(tasks,nr_tasks)pour vous assurer que chaque bloc ne contient que 32 threads et s'inscrit ainsi dans une seule chaîne. Dans cet exemple, j’ai également supposé, par souci de simplicité, que les tâches n’avaient aucune dépendance (par exemple, une tâche dépend des résultats d’un autre) ou des conflits (par exemple, le travail sur la même mémoire globale). Si tel est le cas, le choix de la tâche devient un peu plus compliqué, mais la structure est essentiellement la même.

Ceci est bien sûr plus compliqué que de tout faire sur un seul lot de cellules, mais élargit considérablement le type de problèmes pour lesquels les GPU peuvent être utilisés.

Pedro
la source
2
Cela est techniquement vrai, mais un parallélisme élevé est nécessaire pour obtenir une bande passante mémoire élevée et le nombre d’appels du noyau asynchrone est limité (actuellement, 16). Il existe également des tonnes de comportements non documentés liés à la planification dans la version actuelle. Je vous déconseille de compter sur les noyaux asynchrones pour améliorer les performances pour le moment ...
Max Hutchinson
2
Ce que je décris peut être fait en un seul appel du noyau. Vous pouvez créer N blocs de 32 unités d'exécution chacun, de sorte que chaque bloc s'insère dans une seule chaîne. Chaque bloc acquiert ensuite une tâche à partir d'une liste de tâches globale (accès contrôlé à l'aide d'atomics / mutexes) et la calcule à l'aide de 32 threads verrouillés. Tout cela se passe dans un seul appel du noyau. Si vous souhaitez un exemple de code, faites-le moi savoir et j'en posterai un.
Pedro
4

Un point qui n’a pas été soulevé jusqu’à présent est que la génération actuelle de GPU n’est pas aussi performante pour les calculs à virgule flottante double précision que pour les calculs à simple précision. Si vos calculs doivent être effectués en double précision, vous pouvez vous attendre à ce que le temps d’exécution augmente d’un facteur 10 environ par rapport à une seule précision.

Brian Borchers
la source
Je veux être en désaccord. La plupart (ou tous) les nouveaux GPU ont un support natif en double précision. Presque tous ces GPU signalent des calculs en double précision s’exécutant à une vitesse deux fois moins rapide que celle en simple précision, probablement en raison du simple doublement des accès mémoire / de la bande passante requis.
Godric Seer
1
Il est vrai que les cartes Nvidia Tesla les plus récentes et les plus performantes offrent des performances maximales en double précision qui représentent la moitié des performances maximales en simple précision, mais le rapport est de 8: 1 pour les cartes grand public de l’architecture Fermi les plus courantes.
Brian Borchers
@GodricSeer Le rapport 2: 1 de la virgule flottante SP et DP a très peu à voir avec la bande passante et presque tout avec le nombre d'unités matérielles disponibles pour exécuter ces opérations. Il est courant de réutiliser le fichier de registre pour SP et DP. Par conséquent, l'unité à virgule flottante peut exécuter deux fois les opérations SP en tant qu'opérations DP. Il existe de nombreuses exceptions à cette conception, par exemple, IBM Blue Gene / Q (n’a pas de logique SP et par conséquent, SP s'exécute à environ 1,05x DP). Certains GPU ont des ratios autres que 2, par exemple 3 et 5.
Jeff
Cela fait quatre ans que j'ai écrit cette réponse, et la situation actuelle avec les GPU NVIDIA est que, pour les lignes GeForce et Quadro, le rapport DP / SP est maintenant de 1/32. Les GPU Tesla de NVIDIA offrent des performances beaucoup plus fortes en double précision mais coûtent également beaucoup plus cher. D'un autre côté, AMD n'a pas gêné les performances en double précision de son processeur graphique Radeon de la même manière.
Brian Borchers
4

D'un point de vue métaphorique, le gpu peut être vu comme une personne allongée sur un lit de clous. Les données sont placées au-dessus de la personne. À la base de chaque clou se trouve un processeur. Le clou est donc une flèche pointant du processeur à la mémoire. Tous les ongles sont dans un modèle régulier, comme une grille. Si le corps est bien étendu, il se sent bien (la performance est bonne), si le corps ne touche que quelques points du lit de l'ongle, la douleur est alors mauvaise (mauvaise performance).

Cela peut être considéré comme une réponse complémentaire aux excellentes réponses ci-dessus.

labotsirc
la source
4

Vieille question, mais je pense que cette réponse de 2014 - liée à des méthodes statistiques, mais généralisable pour quiconque sait ce qu'est une boucle - est particulièrement illustrative et informative.

GT.
la source
2

Les GPU ont des E / S de latence longue, il faut donc utiliser beaucoup de threads pour saturer la mémoire. Maintenir une chaîne occupée nécessite beaucoup de threads. Si le chemin de code est de 10 horloges et que la latence des E / S est de 320 horloges, 32 threads devraient presque saturer la chaîne. Si le chemin du code est de 5 horloges, doublez le nombre de threads.

Avec mille cœurs, recherchez des milliers de threads pour utiliser pleinement le GPU.

L'accès à la mémoire se fait par ligne de cache, généralement 32 octets. Le chargement d'un octet a un coût comparable à 32 octets. Donc, coalescez le stockage pour augmenter la localité d'utilisation.

Il y a beaucoup de registres et de RAM locale à chaque chaîne, permettant ainsi le partage entre voisins.

Les simulations de proximité de grands ensembles devraient être bien optimisées.

Les entrées / sorties aléatoires et les threads simples sont une joie mortelle ...

utilisateur14381
la source
C'est une question véritablement fascinante. Je discute avec moi-même pour savoir s'il est possible (ou en vaut la chandelle) de "mettre en parallèle" une tâche relativement simple (détection de bord dans les images aériennes) lorsque chaque tâche prend environ 0,06 s, mais il y a environ 1,8 million de tâches à effectuer ( par an, pour 6 ans de données: les tâches sont clairement séparables) ... soit environ 7,5 jours de temps de calcul sur un cœur. Si chaque calcul était plus rapide sur un GPU et que le travail pouvait être mis en parallèle 1 petit par nGPUcores [n petit], est-il réellement probable que le temps de travail puisse chuter à environ 1 heure? Cela semble peu probable.
GT.
0

Imaginez un problème qui peut être résolu avec beaucoup de force brute, comme Travelling Salesman. Imaginez ensuite que vous avez des racks de serveurs avec chacun 8 cartes vidéo spanky et que chaque carte contient 3 000 cœurs CUDA.

Il suffit de résoudre TOUS les itinéraires possibles du vendeur, puis de trier le temps / la distance / certaines métriques. Bien sûr, vous jetez presque 100% de votre travail, mais la force brute est parfois une solution viable.

Criggie
la source
J'ai eu accès à une petite ferme de 4 serveurs de ce type pendant une semaine et, en cinq jours, j'ai utilisé plus de blocs de type Distributed.net qu'au cours des 10 années précédentes.
Criggie
-1

Après avoir étudié de nombreuses idées d'ingénierie, je dirais qu'un gpu est une forme de focalisation de tâches, de gestion de la mémoire, de calcul répétable.

Beaucoup de formules peuvent être simples à écrire, mais difficiles à calculer, comme dans les mathématiques matricielles, on n'obtient pas une seule réponse mais plusieurs valeurs.

Ceci est important en informatique car la vitesse à laquelle un ordinateur calcule des valeurs et exécute des formules car certaines formules ne peuvent pas s'exécuter sans toutes les valeurs calculées (donc ralentissent). Un ordinateur ne sait pas très bien dans quel ordre exécuter des formules ou calculer des valeurs à utiliser dans ces programmes. Il faut principalement calculer les forces brutes à des vitesses rapides et diviser les formules en mandrins, mais de nombreux programmes exigent de nos jours ces mandrins calculés et attendent en ques (et ques de ques et plus de ques).

Par exemple, dans un jeu de simulation qui devrait être calculé en premier dans les collisions, les dommages de la collision, la position des objets, la nouvelle vitesse? Combien de temps cela devrait-il prendre? Comment un processeur peut-il gérer cette charge? En outre, la plupart des programmes sont très abstraits et nécessitent plus de temps pour traiter les données et ne sont pas toujours conçus pour le multi-threading ou ne constituent pas un bon moyen pour les programmes abstraits de le faire efficacement.

Au fur et à mesure que le processeur devenait de mieux en mieux, la programmation devenait de plus en plus médiocre et nous devions également programmer pour différents types d'ordinateurs. Un gpu est conçu pour forcer la force à travers plusieurs calculs simples en même temps (sans parler de la mémoire (secondaire) ou du bélier) et le refroidissement par chauffage est le principal goulot d'étranglement de l'informatique. Un processeur gère beaucoup de questions en même temps ou est entraîné dans de nombreuses directions, il s'agit de déterminer ce qu'il faut faire si on ne peut pas le faire. (hé c'est presque humain)

Un gpu est grunt travailleur le travail fastidieux. Un processeur gère un chaos complet et ne peut gérer tous les détails.

Alors qu'est-ce qu'on apprend? Un gpu détaille le travail fastidieux en une seule fois et un cpu est une machine multi-tâches qui ne peut pas se concentrer très bien avec trop de tâches à effectuer. (C'est comme s'il avait un trouble de l'attention et l'autisme en même temps).

L'ingénierie comprend les idées, la conception, la réalité et beaucoup de travail ingrat.

Lorsque je pars, souvenez-vous de commencer simplement, commencez rapidement, échouez vite, échouez vite, et n'arrêtez jamais d'essayer.

Andrew G. Corbi
la source