GPU et projets académiques

Séminaire départemental
David Landry
Université Laval
24 février 2017

À propos de moi●●●Étudiant à la maîtrise
en Robotique Mobile
Sous la supervisionde Philippe Giguère
Perception avec
LiDAR– Nuages de points

David Landry

GPU et projets académiques2Objectif
Susciter l’utilisation du GPGPU(General Purpose Graphics Processing Unit)
dans les projets des étudiants du département

David Landry

GPU et projets académiques

3

Résumé(1) Motivation(2) Qu’est-ce qu’un GPU?
(3) Comment détecter un programmeparallélisable?
(4) Où commencer?

David Landry

GPU et projets académiques

4

Résumé(1) Motivation(2) Qu’est-ce qu’un GPU?
(3) Comment détecter un programmeparallélisable?
(4) Où commencer?

David Landry

GPU et projets académiques

5Pourquoi s’y intéresser?
« C’est seulement un petit projet, lesperformances ne sont pas importantes »David Landry

GPU et projets académiques6Cycle de développement
CodeTest
Ça marche pas :(David Landry

GPU et projets académiques7Cycle de développement●●●Plus le cycle est facile, plus on a envie de lerépéter
Plus on répète le cycle, meilleur est le code
Un code plus rapide favorise l’exploration deparamètres, de nouvelles solutions– La confection de réseaux de neurones est un bon

exempleDavid Landry

GPU et projets académiques

8Pourquoi s’y intéresser?
« Je vais perdre plus de temps à apprendre à
programmer des GPU que je vais en sauver »David Landry

GPU et projets académiques

9Pourquoi s’y intéresser?
●●Il existe maintenant des librairies facilesd’utilisation qui permettent une amélioration
rapide du code
On passe toujours plus de temps que prévu surun projet– Le ratio temps investi/temps sauvé est donc

meilleur qu’on le penseDavid Landry

GPU et projets académiques

10

Résumé(1) Motivation(2) Qu’est-ce qu’un GPU?
(3) Comment détecter un programmeparallélisable?
(4) Où commencer?

David Landry

GPU et projets académiques

11GPU: un historique●●Initialement créés pour accélérer le rendu graphique–Opérations massivement parallèles (pour chaque pixel)–Multiplications de matrices~1998 on peut pour la première fois exécuter son
propre code sur des GPU (shaders)–●On soit encore programmer en terme de primitives graphiques (pixels,polygones...)~2006 CUDA permet de programmer sur le matériel
Nvidia en C–Plus besoin de reformuler le problème. On parle de GPGPU.

David Landry

GPU et projets académiques

12Interfaces vers le GPGPU●

OpenCL est l’API libre le plus
commun pour développer surGPU– Il est aussi capable de cibler desCPU multi-coeurs●●David Landry

CUDA est l’api propriétaire
dominant
Dans les deux cas, on écrit ducode pour GPU en C
GPU et projets académiques

13Interfaces vers le GPGPU●●●Dans les deux cas, l’écriture de boncode demande des connaissances
spécifiques
Heureusement, des librairies ont étéconstruites par dessus ces interfaces
bas-niveau
C’est sur les librairies que nous allons
nous concentrer– Même pour utiliser les librairies, il fautêtre conscient du fonctionnement desGPUDavid Landry

GPU et projets académiques

14Interfaces vers le GPGPU●OpenCL est supporté par unetrès grande variétéd’accélerateurs– AMD Radeon– Nvidia– Intel●David Landry

CUDA est supporté seulement
par les appareils Nvidia

GPU et projets académiques

15Comparatif des architectures
Haut-débit

Basse latence

Programmes sur desdonnées en flux

Programmes générauxDavid Landry

GPU et projets académiques

16Comparatif des architectures
Basse latenceHaut-débit

Les opérations individuellessont rapides

David Landry

Fait plusieurs opérations
plus lentement, maisen parallèle

GPU et projets académiques17Interactions avec GPU
Données

KernelRésulat24 go de RAM(Tesla K80)

8-32 go de RAM

David Landry

GPU et projets académiques18Interactions avec GPUDonnées

KernelRésulat8-32 go de RAM

David Landry

Transferts coûteux
GPU et projets académiques

24 go de RAM(Tesla K80)

19

Interactions avec GPU
Morale de l’histoire
Il faut limiter les interactions entre le CPU et le GPU

David Landry

GPU et projets académiques20À garder en têteBranch Divergence●Une carte graphiquecontient plusieurs Streaming
Multiprocessors (SM)– Par exemple une GTX 1080

contient 20 SM●

Chaque SM peut exécuter
entre 1-6 warps à la fois
– Un warp est un groupe de 32threads●Les warps exécutent une
seule instruction à la foisDavid Landry

GPU et projets académiques21À garder en têteBranch Divergence●Comme les warps exécutent une seuleinstruction à la fois, les divergences ont un
impact négatif sur les performances– Par exemple, si un kernel contient un if-then-else,on exécute le then sur tous les threads concernés,
puis le else sur tous les autres threads– L’exécution des deux blocs est faiteséquentiellement

David Landry

GPU et projets académiques22À garder en têteBranch Divergence
Morale de l’histoire
Le GPU ne sera pas utilisé efficacement si votre
kernel contient beaucoup des branchements

David Landry

GPU et projets académiques23À garder en têteBlock Scheduling●●●

Quand on transfère les données au GPU, on
les divise en blocs
Les blocs doivent être indépendants entre eux
et pouvoir s’exécuter dans n’importe quel ordre
Ces conditions permettent au processeur depasser d’un bloc à l’autre sans problèmeDavid Landry

GPU et projets académiques24À garder en têteBlock Scheduling
●●

●

●

Chaque SM peut avoir jusqu’à 8 blocs enréserve
Si un bloc est en attente de données, le SM
active un autre bloc pour continuer les calculs
C’est de cette façon qu’on dissimule la latence
derrière des calculs
Un bon nombre de blocs est donc une bonne
choseDavid Landry

GPU et projets académiques25À garder en têteBlock Scheduling
Morale de l’histoire
Le GPU ne sera pas utilisé efficacement si on n’apas suffisament de donnéesDavid Landry

GPU et projets académiques

26

Résumé(1) Motivation(2) Qu’est-ce qu’un GPU?
(3) Comment détecter un programmeparallélisable?
(4) Où commencer?

David Landry

GPU et projets académiques

27Comment détecter lesapplications du GPU?
●En règle générale, on sait déjà qu’il faut
– Éviter les interactions GPU/CPU fréquentes– Éviter les branchements (if)– Avoir une grande quantité de données à traiter●●

Il faut avoir des données en blocs traitables
indépendamment
Ces règles abstraites sont mieux expliquéesavec des exemples

David Landry

GPU et projets académiques28Collection arbitraire et non-exhaustive destypes de programmes portables sur GPU
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

Données contigües
Itérables dans ledésordre

David Landry

GPU et projets académiques29Traitement d’images
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

●Les GPUs ont été conçus
pour traiter des matrices depixels, donc c’est
nécessairement un bon
exemple d’utilisation– Filtration– Détection de coins

Données contigües
Itérables dans ledésordre

David Landry

GPU et projets académiques

30

Opérations sur les vecteurs
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

●Min, max

●Sum●Sort

●Uniq

●

Dérivation numérique

Données contigües
Itérables dans ledésordre

David Landry

GPU et projets académiques

31

Algèbre linéaire
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

Données contigües
Itérables dans ledésordre

David Landry●●●

Multiplications de matrices
Systèmes d’équations
linéaires
Factorisations– Un des premiers

programmes scientifiques àbénéficier des GPU est la
factorisation LU

GPU et projets académiques

32Données contigües et itérables dans ledésordre
Traitementd’images●Votre code peut probablement êtreporté sur GPU si vos données sont...
– Contigües en mémoire

Opérations surles vecteurs– Itérables dans le désordre●Algèbre linéaire
●Données contigües
Itérables dans ledésordre

David Landry

Des opérations plus compliquées
peuvent l’être aussi, mais c’est ducas par cas
Parfois il faut travailler le problèmede façon à le ramener sous une
forme parallélisable

GPU et projets académiques33

Résumé(1) Motivation(2) Qu’est-ce qu’un GPU?
(3) Comment détecter un programmeparallélisable?
(4) Où commencer?

David Landry

GPU et projets académiques34

Le matériel
●●

●

●

●Les GPU Nvidia sont « dans le vent »Il existe un GPU Grant Program (professeursseulement) http://bit.ly/gpugrant
Un portable avec GPU peut faire une bonne partie duchemin
Le super-ordinateur Helios, de Calcul Québec, a desGPU sur tous les noeuds
Dans le pire des cas, on en trouve pour 100-1000$
sur le marché, selon les performances désiréesDavid Landry

GPU et projets académiques35

Interface de programmation●On cherche à éviter les interfaces bas niveau
comme CUDA et OpenCL– C’est difficile à déboguer– Il faut être très familier avec l’architecture pourécrire du code performant– Le code sera plus difficilement portable à d’autres

cibles d’exécution●On va plutôt préférer utiliser des abstractions
qui utilisent le GPU pour nous

David Landry

GPU et projets académiques36Choisir ses outils
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

Données contigües
Itérables dans ledésordre

David Landry

GPU et projets académiques37

Choisir ses outils
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

Données contigües
Itérables dans ledésordre

David Landry

GPU et projets académiques

38

OpenCV« OpenCV est une bibliothèquegraphique libre spécialisée dans letraitement d’images en temps réel »●Licence: BSD●

Disponible depuis le code source et sur Ubuntu●Utilisable avec Python, C++, JavaDavid Landry

GPU et projets académiques

39

OpenCV
Pour accélérer les calculs de OpenCV avec GPUil faut●●●Installer CUDA
Faire un build depuis le code source avec
quelques flags Cmake (ex. -D
WITH_CUDA=ON)
Indiquer à OpenCV de faire les transferts
CPU/GPU dans les programmesDavid Landry

GPU et projets académiques

40

OpenCV
Visite d’un code d’exemple

David Landry

GPU et projets académiques

41Gains de performanceDavid Landry

GPU et projets académiques42Choisir ses outils
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

Données contigües
Itérables dans ledésordre

David Landry

GPU et projets académiques

43

●Librairie multi-fonctions qui peut s’exécuter surplusieurs cibles– CUDA (GPU Nvidia)– OpenCL (Presque tous les accélérateurs)– CPU (si nécessaire)●

Utilisable en C++, Python, Rust, Node.js– Compatible avec Numpy●

Licence: BSD

David Landry

GPU et projets académiques

44Une visite sur la documentation de la librairie
indique qu’elle a des modules pour des domainesvariés●Algorithmes sur des vecteurs●Algèbre linéaire

●Traitement de signal●StatistiquesDavid Landry

GPU et projets académiques

45

Pour accélérer ses calculs avec ArrayFire il faut●●●Installer CUDA
Se procurer ArrayFire (depuis les paquets
Ubuntu ou depuis le code source)[Optionnel] Installer l’interface Python

David Landry

GPU et projets académiques46

Démonstration

David Landry

GPU et projets académiques

47

Exemples de gains
Opération

Speedup
Nvidia GT630

Produit scalaire

Speedup
Nvidia Tesla
1x

10xMultiplication matrice-vecteur

1.5x

16x

Ordonner un vecteur

14x

80x

0.06x

190x

Convolution 2D

42x

600x

Multiplication matrice-matrice

210x

2000x

Système d’équations linéaires

David Landry

GPU et projets académiques48

Choisir ses outils
Traitementd’images

Opérations surles vecteurs

Algèbre linéaire

Données contigües
Itérables dans ledésordre

David Landry

GPU et projets académiques49« The OpenACC Interface describes acollection of compiler directives to specify
loops and regions of code [...] to be offloaded
from a host CPU to an attached accelerator. »●Licence: Propriétaire●

Utilisable avec C, C++ (et Fortran)

David Landry

GPU et projets académiques

50●●OpenACC est un ensemble de directives quiindiquent au compilateur quelles sections ducode sont parallélisables
De là, le compilateur fait de son mieux pourporter le code sur GPU– Encore une fois, le code qu’on écrit n’est passpécifique à notre cible d’exécution●La code est portable sur des accélérateurs
Nvidia, AMD, et IntelDavid Landry

GPU et projets académiques51

Ceci indique aucompilateur de#pragma acc kernels{

1) Identifier le parallélisme
dans le codefor (int i=0; i < N; i++) {C[i] = A[i] + B[i];}2) Identifier les données àtransférer sur GPU}

3) Créer un kernel
4) Faire exécuter le calcul
sur GPUDavid Landry

GPU et projets académiques52

Pour la programmation GPU par directives, uncompétiteur de OpenACC est OpenMP

Propriétaire

Open source

Principalement descriptif

Principalement prescriptif

David Landry

GPU et projets académiques

53

●

●

●OpenACC est un standard implémenté surplusieurs compilateurs
Le support par le compilateur GCC n’est pas
encore au point
Il faut donc se procurer le compilateur
propriétaire PGI– La Community Edition est disponible gratuitement

David Landry

GPU et projets académiques

54Visite d’un code d’exemple

David Landry

GPU et projets académiques

55

Synthèse – Librairies GPUExpérience GPU requiseDavid Landry

GPU et projets académiques

56

Synthèse – Librairies GPULibrairie

C/C++

JavaPython

X

X

X

XX

X

David Landry

GPU et projets académiques

57

Ressources importantes

Prochaines formations

http://bit.ly/formations2David Landry

GPU et projets académiques

58

Ressources importantes

https://developer.nvidia.com/gpu-accelerated-libraries
David Landry

GPU et projets académiques59

Merci! Des questions?

Merci à Marc Parizeau pour les conseilsDavid Landry

GPU et projets académiques60

La librairie Thrust« Thrust is a parallel algorithms librarywhich resembles the C++ Standard
Template Library (STL). »●Licence: Apache●Disponible gratuitement sur internet

David Landry

GPU et projets académiques61La librairie Thrust●●●David Landry

Essentiellement, elle réimplémente
certaines opérations de la STL sur GPU
Des algorithmes tels que max, sum, sort,partition... sont donc disponibles
facilement
On peut aussi coder ses propres kernels
avec des fonctions telles que transformou reduce (moins recommandé)

GPU et projets académiques

62Comment se laprocurer●●●David Landry

Elle est fournie avec le toolkit CUDA deNvidia
Il faut donc la télécharger et l’installer à
http://bit.ly/get-cuda
Si l’installation est réussie, Thrust va
s’utiliser comme une librairie C++
normale.

GPU et projets académiques

63

ExempleDavid Landry

GPU et projets académiques

64