1 / 90

GPA667: Conception et simulation de circuits électroniques Calcul parallèle avec GPU

GPA667: Conception et simulation de circuits électroniques Calcul parallèle avec GPU. Sommaire. Définition du GPU = ‘ Graphical Processor Unit ’ un engin de traitement graphique très puissant ‘ General Purpose Processing on a GPU’ (GPGPU):

zinna
Télécharger la présentation

GPA667: Conception et simulation de circuits électroniques Calcul parallèle avec GPU

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. GPA667: Conception et simulation de circuits électroniques Calcul parallèle avec GPU GPA667: Conception et simulation de circuits électroniques

  2. Sommaire Définition du GPU = ‘Graphical Processor Unit’ • un engin de traitement graphique très puissant • ‘General PurposeProcessing on a GPU’ (GPGPU): • depuis 2004, le GPU est perçu comme un processeur programmable qui est massivement parallèle • grande capacité de traitement dans chaque PC! • GPU vs CPU: • architecture et modèle du programmeur très différentes • performances nettement plus élevées pour certaines classes d’applications GPA667: Conception et simulation de circuits électroniques

  3. Sommaire – Section E Documents de référence: • Hwu & Kirk, ProgrammingMassivelyParallel Processors, Morgan Kaufmann, 2010 Chapitres du livre sur le site GPA667 • NVIDIA, NVidia CUDA Programming Guide, NVIDIA, version 2.3.1, 145 pages, 2009 GPA667: Conception et simulation de circuits électroniques

  4. Sommaire Calcul parallèle avec GPU: • Notions préliminaires • Intro. à la reconnaissance de formes • Modèle de programmation CUDA • Mémoires et transfert de données • Exécution et ordonnancement de tâches • Évaluation des performances GPA667: Conception et simulation de circuits électroniques

  5. 1 Notions préliminaires • Évolution des microprocesseurs • CPU basé sur un seul ALU: • 1980 – 2003: à l’origine d’une croissance rapide de performances et fonctionnalités pour plusieurs applications • depuis 2003: le limitations de fclk et de consommation de puissance du CPU occasionnent un ralentissement • Tendance actuelle – CPU avec plusieurs ALU: • circuits multi-cœurs pour augmenter les performances • défis: assigner des séquences d’instructions aux cœurs • le microprocesseur est maintenant une machine parallèle! GPA667 Section E: Calcul parallèle avec GPU

  6. 1 Notions préliminaires • Évolution des microprocesseurs • Programmes séquentielles: • les applications sont traditionnellement réalisées avec des programmes de nature séquentiels • ces programmes exécutent une séquence d’instructions sur un seul ALU • Tendance actuelle – réalisations avec programmes parallèles: • ces programmes exécutent plusieurs sous-séquences d’instructions qui coopèrent pour augmenter les performances GPA667 Section E: Calcul parallèle avec GPU

  7. 1 Notions préliminaires • Émergence du GPU • GPU ≡ un classe de processeurs multi-coeurs • Performances GPU vs CPU: • débit des calculs: 367 GFLOPS vs. 32 GFLOPS • bande passante mémoire: 86.4 GB/s vs. 8.4 GB/s G80 = GeForce 8800 GTX G71 = GeForce 7900 GTX G70 = GeForce 7800 GTX NV40 = GeForce 6800 Ultra NV35 = GeForce FX 5950 Ultra NV30 = GeForce FX 5800 GPA667 Section E: Calcul parallèle avec GPU

  8. 1 Notions préliminaires • Émergence du GPU • Architecture CPU: • - optimisé pour exécuter le code séquentiel (1 séquence) • - grande cache avec la bande passante mémoire limitée • Architecture GPU: • - optimisé pour exécuter un grand nombre de séquences • - petites caches avec grande bande passante mémoire GPA667 Section E: Calcul parallèle avec GPU

  9. 1 Notions préliminaires • Architecture d’un GPU GPA667 Section E: Calcul parallèle avec GPU

  10. 1 Notions préliminaires • Architecture d’un GPU • Terminologie (technologie G80): • 16 multiprocesseurs, SM(‘Streaming Multiprocessors’), soit 2 par bloc • 8 cœurs SP (‘Scalar Processor’) par SM • chaque SP comprend des unités fonctionnels: • MAC et MUL roulant à 1.35 GHz • SQRT et RCP SQRT point flottant • Total: 16 SM x 8 SP/SM = 128 cœurs (ou ALU) • 500+ GFLOPS possible GPA667 Section E: Calcul parallèle avec GPU

  11. 1 Notions préliminaires • Architecture d’un GPU • Terminologie (technologie G80): • mémoire: • 1.5 MB DRAM cache avec une grande bande passante mémoire de 86.4 GB/s • bande passante de 4 GB/s par direction sur le bus PCI-express (communications avec le CPU hôte et mémoire) • performance: • maximum: 10+ fois plus de FLOPS qu’un CPU de haute gamme • séquences simultanées: supporte 768 séquences/cœur, environ 12,000 séquences/C.I. GPA667 Section E: Calcul parallèle avec GPU

  12. 1 Notions préliminaires • Applications GPU • Pour une classe particulière d’applications, une réalisation GPU permet une grande accélération • Propriétés: • le besoin de calculs est énorme: on traite beaucoup de données • le parallélisme est considérable: le calcul peut être parallélisé pour différentes parties, et à différents niveaux de granularité • le débit de traitement est élevé GPA667 Section E: Calcul parallèle avec GPU

  13. 1 Notions préliminaires • Applications GPU • L’accélération dépend: • de la proportion d’une application qui peut être parallélisée • …si une grande proportion du temps d’exécution est passée dans la partie parallélisable, l’accélération sera grande • des limitations de la bande passante mémoire • de la capacité de la mémoire cache • Complémentarité: • assurer que le GPU vient supporter l’exécution du CPU • les applications ont des parties qui s’exécutent mieux avec CPU, et d’autres avec GPU GPA667 Section E: Calcul parallèle avec GPU

  14. Sommaire Calcul parallèle avec GPU: • Notions préliminaires • Intro. à la reconnaissance de formes • Modèle de programmation CUDA • Mémoires et transfert de données • Exécution et ordonnancement de tâches • Évaluation des performances GPA667: Conception et simulation de circuits électroniques

  15. Sommaire Calcul parallèle avec GPU: • Notions préliminaires • Intro. à la reconnaissance de formes • Modèle de programmation CUDA • Mémoires et transfert de données • Exécution et ordonnancement de tâches • Évaluation des performances GPA667: Conception et simulation de circuits électroniques

  16. 3 Modèle de programmation CUDA • CUDA (“ComputeUnifiedDeviceArchitecture”) est un modèle de programmation générique • comme extension de C • comme instance d’un modèle de programmation parallèle SPMD (‘Single-Program, Multiple-Data’) • Dans CUDA, un système ordiné comprend • 1 hôte: CPU traditionnel (processeur séquentiel) • 1 ou + composants: GPU (1 ou + processeurs massivement parallèles équipées avec plusieurs unités arithmétiques) GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  17. 3 Modèle de programmation CUDA • Interaction entre hôte (CPU) et composants (GPU): • le code principal qui exécute sur le CPU peut entamer des séquences d’instructions sur une GPU • le GPU agit comme coprocesseur dédié, pour accélérer l’exécution du code parallélisable • i.e., sections du code qui incorporent un bonne proportion d’opérations parallélisables selon des structures de données au programme • parallélisme de données: quand plusieurs opérations arithmétiques peuvent s’exécuter simultanément et indépendamment l’un de l’autre sur les données GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  18. 3 Modèle de programmation CUDA • Exemple de système ordiné selon CUDA: CPU (hôte) GPU avec DRAM locale (composant) GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  19. Texture Texture Texture Texture Texture Texture Texture Texture Texture Host Input Assembler Thread Execution Manager Parallel DataCache Parallel DataCache Parallel DataCache Parallel DataCache Parallel DataCache Parallel DataCache Parallel DataCache Parallel DataCache Load/store Load/store Load/store Load/store Load/store Load/store Global Memory 3 Modèle de programmation CUDA • Mode CUDA G80: • exécute plusieurs séquence • mode d’opération avec interface de programmation GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  20. 3 Modèle de programmation CUDA • Exemple de parallélisme de données: multiplication matricielle PIxJ = MIxKxNKxJ …chaque élément pij de la matrice P est obtenu en calculant le produit scalaire entre la rangée i de M et la colonne j de N. GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  21. 3 Modèle de programmation CUDA • Exemple de parallélisme de données: // Multiplication matricielle: code C sur CPU en précision double void MatrixMulOnHost(float* M, float* N, float* P, int Width)‏ { for (int i = 0; i < Width; ++i)‏ for (int j = 0; j < Width; ++j) { double sum = 0; for (int k = 0; k < Width; ++k) { double a = M[i * width + k]; double b = N[k * width + j]; sum += a * b; } P[i * Width + j] = sum; } } GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  22. 3 Modèle de programmation CUDA Réalisations CUDA • Processus de conception avec CUDA: • identification des parties du programme à paralléliser • isoler les données requises par les parties parallélisées, et allocation de la mémoire sur GPU (fonction API) • transfert des données au GPU (fonction API) • développer un fonction noyau (‘kernel’) qui sera exécuté par les séquences individuelles dans les parties parallèles API ≡ ‘Application Programming Interface’ fonctions donnant accès au GPU GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  23. 3 Modèle de programmation CUDA Réalisations CUDA • Programme CUDA: un code source dont les parties sont soit exécutées sur CPU ou sur GPU • le compilateur NVIDIA C (NVCC) sépare les parties • code CPU: • correspond aux parties avec aucun ou peu de parallélisme • réalisé avec un code ANSI C • compilé avec un compilateur standard • exécuté comme un processus normal sur le CPU hôte GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  24. 3 Modèle de programmation CUDA Réalisations CUDA • Programme CUDA: un code source dont les parties sont soit exécutées sur CPU ou sur GPU • code GPU: • correspond aux parties avec beaucoup de parallélisme • réalisé avec un code GPU (code ANSI C étendu) • la syntaxe associée aux fonctions noyaux indique les parties parallélisées et ainsi que leurs structures de données • compilé avec le compilateur NVCC • lorsqu’un noyau un appelé, il est exécuté sur le GPU GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  25. 3 Modèle de programmation CUDA Réalisations CUDA • Fonction noyau (‘kernel’) : • génère plusieurs séquences afin d’exécuter une partie parallèle → pour tirer profit du parallélisme de données • grille ≡ ensemble de séquences généré lors de l’appel d’une fonction noyau • les noyaux CUDA sont très légers: • CPU (multi-cœur): prennent des milliers de cycles pour générer et céduler quelques séquences • CUDA: un support matériel efficace permet générer et céduler avec peu de cycles des 1000s de séquences GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  26. 3 Modèle de programmation CUDA Réalisations CUDA • Exécution d’un programme CUDA typique: GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  27. 3 Modèle de programmation CUDA Réalisations CUDA • Exemple: fonction main pour la multiplication matricielle GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  28. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • CPU et GPU ont des espaces mémoires propres • les GPU sont réalisés sur cartes électroniques avec leur propre DRAM • Avant d’exécuter un noyau sur GPU… • allocation de mémoire en GPU • transfert de données du CPU à cette mémoire en GPU • Après l’exécution du noyau… • transfert des résultats du GPU en mémoire CPU • libérer la mémoire GPU … les fonctions API CUDA: extensions du langage ANSI C appelés au code pour gérer les mémoires et transferts GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  29. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Modèle de mémoire du GPU (composant CUDA) GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  30. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Mémoire globale: (notre focus pour l’instant) • moyen principal pour communiquer (R/W) des données entre le CPU et le GPU • le contenu est visible par toutes les séquences Grid Block (0, 0)‏ Block (1, 0)‏ Shared Memory Shared Memory Registers Registers Registers Registers Thread (0, 0)‏ Thread (1, 0)‏ Thread (0, 0)‏ Thread (1, 0)‏ Host Global Memory GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  31. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Fonctions API CUDA pour la gestion de la mémoire globale GPU GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  32. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Exemple: • allocation d’une matrice 64x64 à précision simple (float de 32 bits) • attacher la structure à Md WIDTH = 64; Float* Md; int size = WIDTH * WIDTH * sizeof(float); cudaMalloc((void**)&Md, size); cudaFree(Md); GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  33. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Fonctions API CUDA pour le transfert de données entre mémoires GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  34. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Exemple: • transfert au GPU d’une matrice 64x64 à précision simple • la mémoire CPU est M et mémoire GPU est Md • cudaMemcpyHostToDevice et cudaMemcpyDeviceToHost sont des constantes symboliques cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost); GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  35. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Exemple: fonction MatrixMulOnDevice() void MatrixMulOnDevice(float* M, float* N, float* P, int Width)‏ { int size = Width * Width * sizeof(float); float* Md, Nd, Pd; … 1. // Allocate and Load M, N to device memory cudaMalloc(&Md, size); cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); cudaMalloc(&Nd, size); cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice); // Allocate P on the device cudaMalloc(&Pd, size); GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  36. 3 Modèle de programmation CUDA Mémoires GPU et transfert de données • Exemple: fonction MatrixMulOnDevice() 2. // Kernel invocation code – to be shown later … 3. // Read P from the device cudaMemcpy(P, Pd, size,cudaMemcpyDeviceToHost); // Free device matrices cudaFree(Md); cudaFree(Nd); cudaFree (Pd); } GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  37. 3 Modèle de programmation CUDA Fonctions noyaux • Appel d’une fonction noyau CUDA: • indique au compilateur NVCC le code qui doit être exécuté en parallèle • appelé par une fonction CPU pour générer des grilles de plusieurs séquences d’instructions • tous les séquences exécutent le même code noyau GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  38. 3 Modèle de programmation CUDA Fonctions noyaux • Extensions de la syntaxe ANSI C: • __global__nom_fonction: indique au compilateur NVCC que la fonction est de type noyau • blocIdx.xetblocIdx.y: • indices uniques permettant de distinguer des blocs de séquences dans une grille • threadIdx.x, threadIdx.yetthreadIdx.z: • indices uniques permettant de distinguer les séquences dans un bloc • permet à une séquence d’accéder ses registres en temps-réel GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  39. Exécuté sur: Appelé sur: __device__ float DeviceFunc()‏ GPU GPU __global__ void KernelFunc()‏ GPU CPU __host__ float HostFunc()‏ CPU CPU 3 Modèle de programmation CUDA Fonctions noyaux • Extensions de la syntaxe ANSI C: • __global__: définit un fonction noyau • doit retourner void GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  40. threadID 0 1 2 3 4 5 6 7 … float x = input[threadID]; float y = func(x); output[threadID] = y; … 3 Modèle de programmation CUDA Fonctions noyaux • Appel d’une fonction noyau CUDA: • chaque séquence exécute le même code (SPMD) • chaque séquence a une identité qu’il utilise pour calculé des adresses mémoire et pour le contrôle GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  41. 3 Modèle de programmation CUDA Fonctions noyaux • Exemple: fonction MatrixMulOnDevice() GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  42. 3 Modèle de programmation CUDA Nd Fonctions noyaux • Exemple: for (int k = 0; k < Width; ++k)‏ { float Melement = Md[threadIdx.y*Width+k]; float Nelement = Nd[k*Width+threadIdx.x]; Pvalue += Melement * Nelement; } Pd[threadIdx.y*Width+threadIdx.x] = Pvalue; } k WIDTH tx Md Pd ty ty WIDTH tx k 42 WIDTH WIDTH GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  43. 3 Modèle de programmation CUDA Fonctions noyaux • Une fonction noyau est exécutée avec une grille de séquences parallèles: • grille: typiquement environ 103 séquences par appel pour pleinement exploiter le GPU • les séquences d’une grille sont organisées selon un hiérarchie à deux niveaux • chaque grille comprend 1+ blocs avec le même nombre de séquences (blockIdx.x, y) • chaque bloc est organisé de la même façon, comme un matrice 3D de séquences(threadIdx.x, y, z) • Max: 512 séquences / bloc GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  44. 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 threadID … float x = input[threadID]; float y = func(x); output[threadID] = y; … … float x = input[threadID]; float y = func(x); output[threadID] = y; … … float x = input[threadID]; float y = func(x); output[threadID] = y; … 3 Modèle de programmation CUDA Fonctions noyaux • Organisation des grilles en blocs multiples: • les séquences dans un bloc peuvent coopérer: mémoire partagée, opérations atomiques et synchronisation • les séquences de différents blocs ne peuvent coopérer Thread Block 1 Thread Block N - 1 Thread Block 0 … GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  45. 3 Modèle de programmation CUDA Fonctions noyaux • Un noyau est exécuté comme une grille de séquences parallèles: GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  46. 3 Modèle de programmation CUDA Fonctions noyaux • Identification des blocs et séquences: Chaque bloc et séquence sont identifiées • bloc: 1D ou 2D • séquence: 1D, 2D ou 3D afin de simplifier l’adressage mémoire Traitement de données multi- dimensionnels (e.g., images) GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  47. 3 Modèle de programmation CUDA Fonctions noyaux • Exemple d’appel d’un noyau par CPU: __global__ void KernelFunct(...); dim3 DimGrid(100, 50); // 5000 thread blocks dim3 DimBlock(4, 8, 8); // 256 threads per block size_t SharedMemBytes = 64; // 64 bytes of shared memory KernelFunct<<< DimGrid, DimBlock, SharedMemBytes >>>(...); GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  48. 3 Modèle de programmation CUDA Fonctions noyaux • Appel d’un noyau par CPU: • Fixer les dimensions du bloc et de la grille en passant de paramètres fonction_noyau‹‹‹dimGrid, dimBlock,SharedMemBytes>>>(paramètres…) GPA667 Section E: Calcul parallèle avec GPU Éric Granger

  49. Sommaire Calcul parallèle avec GPU: • Notions préliminaires • Intro. à la reconnaissance de formes • Modèle de programmation CUDA • Mémoires et transfert de données • Séquences d’instructions CUDA • Évaluation des performances GPA667: Conception et simulation de circuits électroniques

  50. Grid Block (0, 0) Block (1, 0) Shared Memory Shared Memory Registers Registers Registers Registers Thread (0, 0) Thread (1, 0) Thread (0, 0) Thread (1, 0) Host Global Memory Constant Memory 4 Mémoires CUDA Réalisation pour la famille G80 • Chaque séquence peut: • très rapidement… • R/W en registres, par séquence • R/W en mém. locale, par séquence • R/W en mém partagée, par bloc • lentement... • R/W en mém globale, par grille • R only en mém constante, par grille (efficace que globale à cause de cache) GPA667 Section E: Calcul parallèle avec GPU Éric Granger

More Related