Détecter du mouvement en temps réel avec CUDA
Développement d'un plugin GStreamer d'accélération GPU en CUDA pour détecter en temps réel des objets en mouvement dans un flux vidéo.
Introduction
Ce projet a été réalisé dans le cadre du cours "GPGPU" de la majeure IMAGE à l'EPITA, et a pour objectif de développer un plugin GStreamer capable d'effectuer une séparation fond/objets mobiles (background subtraction) sur un flux vidéo en temps réel.
L'enjeu majeur était d'atteindre le temps réel (plus de 30 FPS, et idéalement saturer le flux vidéo d'entrée) en déportant l'intégralité du calcul intensif sur le GPU à l'aide de NVIDIA CUDA.
Quelle est la différence entre un traitement CPU et un traitement GPU ?
Je ne sais pas si cette partie est vraiment nécessaire, mais il me semble tout de même important de rappeler la différence entre un traitement CPU et un traitement GPU.
Concrètement, tous les codes classiques, que ce soit Python, C, C++, Java, etc., sont exécutés sur le processeur central (CPU). Ce dernier est très polyvalent et rapide dans la plupart des cas, mais il est limité à quelques cœurs de calcul (souvent 4 à 16 cœurs pour un CPU grand public).
Ceci implique que lors de gros traitements de données (et c'est typiquement le cas pour les opérations d'analyse de vidéos !), le CPU va devoir exécuter les instructions de manière séquentielle, ce qui peut ralentir le processus de manière significative.
Concrètement, c'est assez facile à visualiser : si vous avez un flux vidéo de 30 images par seconde, et que votre CPU met 100 ms pour traiter une seule image, vous ne pourrez jamais atteindre le temps réel, car il faudra 3 secondes pour traiter 30 images.
Une bonne alternative est de transformer ce code CPU en code GPU, qui sera exécuté sur la carte graphique. Le GPU est composé de centaines, voire de milliers de cœurs de calcul, ce qui permet d'exécuter des milliers d'instructions en parallèle. C'est exactement ce dont nous avons besoin pour traiter un flux vidéo en temps réel.
Si on reprend l'exemple précédent, si le traitement séquentiel de tous les pixels d'une image prend 100 ms sur CPU, mais que le GPU peut traiter des milliers de pixels en parallèle, alors le temps de traitement par image chute à quelques millisecondes, permettant d'atteindre le temps réel sans problème.
Mais pourquoi ne pas toujours coder sur GPU alors ?
C'est une très bonne question, et c'est principalement pour deux raisons assez simples :
- Le code GPU est beaucoup plus complexe à écrire que du code sur CPU, et par conséquent beaucoup plus long à développer. Il faut donc être sûr que le gain de performance en vaut la peine.
- Le GPU est un processeur spécialisé dans le calcul parallèle, mais il est beaucoup moins polyvalent que le CPU. Il n'est pas capable d'exécuter tous les types de calculs, et certaines opérations sont beaucoup plus lentes sur GPU que sur CPU. Il faut donc être sûr que le code que l'on souhaite paralléliser est adapté à l'architecture GPU.
Dans notre cas précis, le traitement vidéo est un candidat idéal pour le GPU, car il s'agit d'un flux de données massives (des millions de pixels par seconde) et les opérations à effectuer sur chaque pixel sont relativement simples et indépendantes les unes des autres.
Après cette petite introduction, nous allons maintenant détailler le pipeline algorithmique que nous avons mis en place pour détecter les objets en mouvement dans un flux vidéo en temps réel.
Le Pipeline Algorithmique
Le pipeline en lui-même n'est pas la partie la plus intéressante du projet, donc voici une présentation assez succincte de celui-ci.
Pour isoler les éléments en mouvement, le flux vidéo passe par un pipeline de traitement d'images composé de 6 étapes successives:
- Modélisation de l'arrière-plan à partir des frames précédentes
- Détection de mouvement brut, mais bruité, par comparaison avec le modèle d'arrière-plan
- Nettoyage du masque de mouvement (ouverture morphologique) pour éliminer les faux positifs
- Seuillage par hystérésis pour lisser les contours des objets détectés
- Colorisation des pixels en mouvement sur la vidéo d'origine
- Affichage du résultat final
L'implémentation CPU de référence
Pour savoir si les optimisations GPU valent le coup, la première étape a été de développer une version séquentielle en C++ pur tournant sur le CPU. (c'est peut-être bête, mais si le code tourne déjà en temps réel sur CPU, il n'y a pas besoin de passer sur GPU !)
Cette version de référence nous sert de "vérité terrain" pour valider notre pipeline, pour pouvoir la comparer en termes de performances et de résultats avec la version GPU.
Côté performances, sans surprise, le CPU montre vite ses limites, puisqu'on tourne en moyenne à 5.29 FPS (soit environ 616 secondes pour traiter notre jeu de données de test). C'est très loin du temps réel visé, ce qui valide complètement l'utilité de passer sur GPU !
L'implémentation naïve en CUDA
La transition vers le GPU commence par une adaptation directe de notre logique séquentielle CPU sous forme de threads CUDA.
CUDA est un langage de programmation parallèle développé par NVIDIA (ceux qui créent des cartes graphiques), qui permet d'écrire des programmes exécutables sur les cartes graphiques NVIDIA. Il est basé sur le langage C/C++, mais avec des extensions pour gérer la parallélisation et la gestion de la mémoire GPU.
Pour cette implémentation naïve, on convertit directement nos boucles for imbriquées en grilles de threads 2D (avec des blocs de taille standard (16 \times 16)). Chaque thread se voit attribuer le calcul d'un seul pixel de l'image. Bien qu'il n'y ait encore aucune optimisation matérielle spécifique, cette simple parallélisation massive nous permet de faire passer le débit à 48.92 FPS (un speedup de x9.24). L'objectif du temps réel est techniquement atteint, mais on sent qu'on peut faire beaucoup mieux. (genre beaucoup beaucoup mieux)
Optimisation 1 : L'allocation paresseuse (Lazy Initialization)
En analysant notre version naïve, on remarque que les transferts mémoire entre le processeur et la carte graphique sabotent nos performances.
Dans une boucle vidéo, allouer et libérer des buffers GPU à chaque frame est juste une abomination ! On a donc mis en place une allocation paresseuse, qui fait que tous les buffers nécessaires au pipeline GPU sont alloués une seule et unique fois lors du chargement de la première image. Ainsi, on réduit le trafic mémoire à seulement deux transferts par frame : l'image brute du CPU vers le GPU (HostToDevice) et l'image traitée du GPU vers le CPU (DeviceToHost). Cette simple correction fait grimper le débit à 95.43 FPS !
Optimisation 2 : Types de données et Remplacement de cuRAND
Pour aller chercher encore plus de performance, on a dégainé les profileurs Nvidia Nsight Compute et Nvidia Nsight Systems pour analyser le pipeline et identifier les goulots d'étranglement.
Ces logiciels permettent de visualiser le temps d'exécution de chaque kernel CUDA, ainsi que l'utilisation des ressources GPU (mémoire, occupation des SM, etc.). Grâce à ces outils, on a pu identifier les parties du code qui consommaient le plus de temps et les optimiser.
Le profileur nous a rapidement alerté sur un problème majeur : nous utilisions des variables de type double (double précision), extrêmement lourdes à calculer pour la plupart des GPU grand public. Le passage au type float (simple précision) et le remplacement des fonctions mathématiques comme round par lroundf ont donné un second souffle aux calculs.
Mais le plus gros goulet d'étranglement venait de la bibliothèque de génération de nombres aléatoires cuRAND. Initialiser un état curandState de 48 octets pour chaque pixel consomme un temps fou au démarrage et gaspille une mémoire VRAM colossale (environ 95 Mo pour une simple frame Full HD 1080p).
Pour y remédier, on a écrit notre propre générateur de nombres aléatoires ultra-léger basé sur un algorithme LCG (Linear Congruential Generator) calculé à la volée. Il utilise les coordonnées du pixel et le numéro de la frame comme germe (seed). Résultat : la mémoire VRAM est libérée et le débit monte à 98.96 FPS.
__device__ unsigned int fast_rand(unsigned int seed) {
seed = (seed ^ 61) ^ (seed >> 16);
seed *= 9;
seed = seed ^ (seed >> 4);
seed *= 0x27d4eb2d;
seed = seed ^ (seed >> 15);
return seed;
}
Optimisation 3 : L'hystérésis en Mémoire Partagée (Shared Memory)
La phase de seuillage par hystérésis consiste à propager l'état d'un pixel fort à ses voisins. Initialement, cela se faisait via une boucle CPU/GPU qui copiait les données dans les deux sens à chaque itération.
Cette synchronisation permanente forçait le CPU et le GPU à s'attendre mutuellement, créant une timeline Nsight catastrophique.
Pour corriger cela, on a migré le calcul de l'hystérésis dans la mémoire partagée (Shared Memory) du GPU. En regroupant les calculs par tuiles de (16 \times 16) pixels avec un halo de sécurité de 1 pixel pour les bords, les threads d'un même bloc peuvent collaborer et s'échanger les informations de propagation à la vitesse de la SRAM locale, sans aucun transfert PCIe. Les performances s'envolent alors à 124.34 FPS !
Optimisation 4 : L'ouverture morphologique tuilée
On a appliqué la même stratégie de tuilage en mémoire partagée pour l'étape d'ouverture morphologique (érosion et dilatation). Je ne vais pas trop m'étendre sur le sujet car c'est littéralement la même idée que pour l'hystérésis, mais avec un halo de 2 pixels pour gérer les bords de la tuile.
Après cette étape, le pipeline atteint désormais 127.57 FPS.
Optimisation 5 : Géométrie des blocs de threads
La dernière optimisation a consisté à ajuster la structure de nos blocs de threads pour saturer au mieux les processeurs de streaming (SM) de notre GPU.
En testant plusieurs géométries sous Nsight Compute ((16 \times 16), (8 \times 8), (32 \times 32), et (32 \times 8)), nous avons constaté que la configuration en blocs rectangulaires de (32 \times 8) threads offrait le meilleur compromis d'occupation et d'accès mémoire alignés (coalescence). Ce dernier ajustement nous permet de gratter encore quelques images par seconde pour culminer au score final de 129.51 FPS.
Bilan Global des Performances
Ce projet a été très fastidieux, notamment à cause du fait qu'il a fallu profiler et optimiser progressivement le code à travers les outils NVIDIA Nsight.
Mais, cela nous a permis de transformer un code CPU séquentiel lent en un code beaucoup plus rapide et efficace sur GPU, capable de traiter un flux vidéo Full HD en temps réel !
Voici l'évolution globale du débit de traitement tout au long du projet :
- C++ Référence (CPU) : 5.29 FPS (1.00x)
- CUDA naïve : 48.92 FPS (9.24x)
- CUDA Optimisée (Lazy Alloc) : 95.43 FPS (18.03x)
- CUDA Fast Random (LCG) : 98.96 FPS (18.70x)
- CUDA Hysteresis (Shared Memory) : 124.34 FPS (23.49x)
- CUDA Tiling Opening (Shared Memory) : 127.57 FPS (24.10x)
- CUDA 32x8 Blocs (Version Finale) : 129.51 FPS (24.47x)
Je suis très satisfait de ce projet, car il m'a permis de m'initier à la programmation GPU avec CUDA, et de comprendre les subtilités de l'optimisation des performances sur ce type d'architecture.