Implémentation d'un CNN - LeNet-5 sur GPU
Réalisé par : Elisa DELHOMME & Pierre CHOUTEAU
Si vous voulez continuer ce projet, ou bien utiliser une partie de celui-ci, vous n'aurez pas besoin de grand chose. Le langage utilisé pour ce projet est du Cuda, il vous faudra donc a minima un PC avec une carte graphique Nvidia. Sinon, vous ne pourrez pas utiliser les fonctions qui utilisent le GPU.
Aujourd'hui le langage Cuda n'est encore présent sur aucun IDE, mais comme la compilation et l'exécution se font via la console, il est possible d'utiliser n'importe quel IDE.
Un IDE comprenant la coloration synthaxique du C ou du C++ fait largement l'affaire. Choisissez donc celui qui vous voulez (Jupyter-Lab, VsCode ou encore Sublime Text font largement l'affaire)
Pour compiler un code Cuda, il vous suffit de lancer la commande :
nvcc nomdufichier.cu -o nomdufichier
Quand vous aurez fait cela, vous verrez apparaître un fichier portant le nom "nomdufichier". Vous n'avez donc plus qu'à l'exécuter, et là encore, rien de plus simple. Lancer simplement la commande :
./nomdufichier
PS: Pour que ces commandes fonctionnent il faut bien sûr que vous soyez dans votre dossier de travail. Vous pouvez vous déplacer facilement dans les dossiers grâce à la commande cd.
Les objectifs de ce projet sont :
- Apprendre à utiliser CUDA
- Etudier la complexité d'algorithmes et l'accélération obtenue sur GPU par rapport à une exécution sur CPU
- Observer les limites de l'utilisation d'un GPU
- Implémenter "from scratch" un CNN : juste la partie inférence et non l'entrainement
- Exporter des données depuis un notebook python et les réimporter dans un projet cuda
- Faire un suivi de votre projet et du versionning grâce à l'outil git
A terme, l'objectif final est d'implémenter l'inférence d'un CNN très classique : LeNet-5
La lecture de cet article apporte les informations nécessaires pour comprendre ce réseau de neurones.
https://www.datasciencecentral.com/profiles/blogs/lenet-5-a-classic-cnn-architecture
Dans tout le projet, que ce soit sur CPU ou GPU, on souhaite représenter les matrices sous forme de listes constituées des lignes de la matrice.
Dans la fonction d'initialisation, on initialise la matrice NxP avec des valeurs aléatoires.
La création d'une fonction d’affichage de matrice sous sa forme classique, avec chaque ligne affichée l’une en dessous l’autre, est utile, voire même nécessaire afin de pouvoir vérifier le bon fonctionnement des opérations traitées ou prendre connaissance d'un résultat.
Sur CPU, on additionne deux matrices simplement comme à notre habitude en sommant les coefficients de chacune deux à deux puisque la représentation sous forme de liste n'est pas un frein à cette addition classique.
Ci-dessous un exemple de réalisation sur une addition de matrices 3x3:
Sur GPU, le calcul se base également sur la somme des coefficients deux à deux mais est un peu plus complexe à mettre en oeuvre en raison de la parallélisation des calculs sur GPU.
En particulier, la définition des indices des coefficients matriciels se fait cette fois via les variables dim3 définissant les threads afin de paralléliser (et donc accélérer) les calculs).
On définira par exemple l'indice de la ligne à considérer (au début de la fonction, et qui changera à chaque itération) par:
int lig = blockIdx.y * blockDim.y + threadIdx.y
où blockIdx.y désigne le numéro de la ligne du block dans la grid,
blockDim.y donne le nombre de lignes total de ce block,
threadIdx.y donne le numéro de la ligne du thread appartenant au block.
Ce qui suit est nécessaire pour tout calcul effectué sur le GPU, et est donc également nécessaire pour la réalisation de la multiplication.
Il est, en outre, nécessaire de définir la fonction avec global afin d’effectuer les calculs sur le GPU en l’appelant depuis le CPU, par exemple dans le cas de l'addition de matrices sur GPU:
__global__ void cudaMatrixAdd(float *M1, float *M2, float *Mout, int n, int p);
Ce sera le cas pour toutes les fonctions effectuant des calculs sur le GPU tout en étant appelé depuis le CPU.
Les mémoires des matrices doivent être allouées sur le GPU avec cudaMalloc et copiées depuis le CPU vers le GPU avec cudaMemcpy
Pour que la parallélisation soit fonctionnelle et efficace il est nécessaire de définir en dehors de la fonction (dans le main par exemple) les dimensions des variables dim3.
dim3 block_size(n, m);
dim3 grid_size(1, 1);
où n et m sont les dimensions de la matrice résultante du calcul.
Enfin, la fonction d'addition sur le GPU est appelée via la commande suivante:
cudaMatrixAdd<<<grid_size, block_size>>>(d_M1, d_M2, d_Mout, n, p);
Ci-dessous un exemple de réalisation:
La multiplication de deux matrices sur le CPU se fait de façon habituelle. La seule difficulté réside dans l'indexage correct des coefficients recherchés, les matrices étant sous forme de liste.
Ci-dessous un exemple de réalisation sur une multiplication de matrices 3x3:
Comme pour l'addition, la multiplication sur GPU repose sur le même principe que la multiplication classique mais les indexes lignes et colonnes désirées doivent être définies (exactement selon la même formulation que pour l'addition) avec les variables définissant les threads et blocks.
Ci-dessous un exemple de réalisation sur une multiplication de matrices 3x3:
L'architecture du réseau LeNet-5 est composée de plusieurs couches :
- Layer 1- Couche d'entrée de taille 32x32 correspondant à la taille des images de la base de données MNIST
- Layer 2- Convolution avec 6 noyaux de convolution de taille 5x5. La taille résultante est donc de 6x28x28.
- Layer 3- Sous-échantillonnage d'un facteur 2. La taille résultante des données est donc de 6x14x14.
La génération des données consiste en la création des matrices suivantes sous la forme de tableaux à une dimension:
- la matrice d'entrée dans le réseau 32x32 raw_data initialisée avec des valeurs aléatoires entre 0 et 1.
- la matrice 6x28x28 C1_data résultante de la convolution 2D initialisée avec des valeurs nulles.
- la matrice 6x14x14 S1_data résultante du sous-échantillonnage initialisée avec des valeurs nulles.
- le kernel 6x5x5 C1_kernel permettant la convolution de la layer 1 et initialisé entre 0 et 1.
La convolution se fait exclusivement sur le GPU. De façon analogue à la multiplication, on fait glisser un kernel C1_kernel sur la totalité de la matrice raw_data pour obtenir la matrice résultante C1_data.
Il est également nécessaire de prendre un compte le nombre de kernel (la profondeur de C1_kernel) sur les calculs de convolution afin d'obtenir le nombre de features maps souhaités.
Le sous-échantillonage se fait par une fonction de MeanPooling, à savoir un moyennage sur une fenêtre glissante 2x2 (afin de réduire par 2 les dimensions de raw_data et d'obtenir S1_data).
Celui-ci se fait également sur le GPU depuis un appel du CPU.
Dans l'objectif de parfaire le réseau de neurones, une couche d'activation est requise. Comme on peut le remarquer dans l'article sur l'architecture de LeNet-5, la fonction d'activation utilisée est une tangente hyperbolique (tanh). Celle-ci interviendra après chaque layer de Conv2D.
Afin de se laisser la possibilité d'appeler cette fonction d'activation depuis chaque kernel sur le GPU, on définit cette fois la fonction avec le specifier __device__, et non __global__ pour effectuer les calculs sur le GPU depuis un appel du GPU.
La couche de fonction d'activation retourne une matrice de même dimension que celle qui lui est fournie.
Voci ci-dessous un exemple de réalisation d'une convolution d'une matrice 8x8 par un kernel 5x5 suivie éventuellement d'une fonction d'activation, puis d'un MeanPooling:
Dans cet exemple, on choisit volontairement des matrices aux coefficients simples afin de pouvoir confirmer le bon déroulement des calculs: le kernel nul avec un 2 central permet notamment la vérification rapide des calculs sur une matrice unitaire.
Dans cette dernière partie, on utilise le notebook Python comme référence afin de finaliser notre réseau LeNet5. En particulier, celui-ci nous servira, grâce à un entraînement rapide, d'obtenir les valeurs optimales des poids de chaque couche afin de pouvoir initialiser les kernels de convolution et les poids des couches fully connected de façon à obtenir les meilleurs résultats.
En effet, dans ce projet, on ne désire pas créer la fonction d'entrainement du réseau de neurones en Cuda, car cela est beaucoup plus complexe, et nous aurait pris trop de temps à mettre en place (Descente de gradient, BackPropagation...). Ceci pourrait donc être un bon point de départ pour continuer ce projet :
- Créer l'optimizer
- La fonction de loss
- L'entrainement sur une base de test et de validation.
On construit le réseau en ajoutant couches de convolution et de MeanPooling. Il est également nécessaire de créer une couche de Dense effectuant l'opération W.x + b où W sont les poids et b, les biais appliqués à l'image d'entrée x. Cette fonction fait intervenir les fonctions de multiplication et d'addition sur le GPU.
En outre, afin de prévoir les cas où les matrices W et x ne sont pas carrées, on se propose d'introduire une nouvelle fonction de multiplication (nommée cudaMatrixMultGeneral), basée sur le même principe que celle créée plus haut sur GPU, mais effectuant la multiplication d'une matrice NxP par une matrice PxM pour donner une matrice résultante NxM (celle-ci se trouve dans le fichier Partie3.cu).
Après avoir entraîné le réseau dans le notebook et récupéré les poids et biais de chaque couche, on les utilise pour initialiser les kernels.
Le réseau LeNet5 est désormais entièrement fonctionnel pour une image d'entrée de dimensions 32x32.
Cette partie n'est pas encore complètement terminée. La création du réseau est faite et fonctionne au vu des différents tests, mais la liaison avec les poids pas encore... Il reste à finir cette partie à finir pour que le réseau puisse être totalement fonctionnel.