Le but de notre projet est de créer un réseau de CNN de type LeNet-5. Ce réseau est constitué de deux alternances de CNN et d'average pooling, suivies de deux couches dense.
Au cours de notre projet nous nous proposons de :
- Exploiter un logiciel de gestion de versions décentralisé tel que git.
- Appréhender la programmation en CUDA
- Faire une étude comparative des performances entre nos fonctions GPU /CPU
- Créer à partir d'un langage bas niveau des couches de neurones de type convolutive ou fully connected
- Exploiter les poids d'un modèle entrainé grâce à Keras de facons à paramétrer nos couches convolutives codées en C.
Dans un premier temps nous allons développer des fonctions simples exécutées par le GPU permettant de créer et observer des matrices.
Chaque matrice est un tableau indexé par un pointeur et la fonction matrixinit() permet de créer un tableau de taille NxP remplis de coefficients aléatoires compris entre -1 et 1.
Par la suite nous avons aussi créé une fonction MatrixPrint permettant d'afficher dans la console une matrice de taille n lignes p colonnes.
Cette fonction est primordiale pour juger de la qualité de notre travail car elle permet de visualiser les résultats de nos diverses opérations ainsi que de juger de la justesse des résultats obtenus.
Pour additionner deux matrices sur CPU, il suffit simplement d'ajouter les uns après les autres les coefficients de même indice des deux matrices d'entrée puis de venir stocker ces résultats dans la matrice Mout. Il n'y a ici aucune parallélisation des calculs, les additions et les affectations sont faites par le CPU les unes après les autres. On commence à comprendre que ces calculs pourraient être beaucoup plus rapides si le calcul d'un coefficient et une affectation étaient faits par un thread différent. On gagnerai ainsi un temps proportionnel au nombre de coéfficient.
On obtient cependant un résultat tout à fait correct :
On se propose maintenant de réaliser une fonction cudaMatrixAdd() permettant d'exploiter les capacités de parallélisation de nos GPU. Pour cela il est important de réflechir aux dimensions que nous allons donner à notre Grid ainsi qu'a nos blocs. Nous avons choisi de raisonner de la facons suivante :
- Chaque block correspond à une une ligne des matrices
- Chaque thread au sein de ces blocks correspondent à une colonnes des matrices.
En définissant la dimension de la facons suivante :
On en déduit la fonction suivante :
On observe que cette fonction permet à chaque thread de chaque block de calculer un coefficient différent de l'adition de matrice et de le stocker au bon endroit dans la matrice de sortie Mout.
En cuda les matrices sont indicées grâce aux variables :
- gridDim.x/y représentant la dimension totale selon x/y de la grille que nous avons défini.
- blockDim.x/y représentant la dimension totale selon x/y de chaque block que nous avons défini au sein de la grid.
- blockIdx.x/y représentant l'Id du block (selon x ou y) auquel appartient le thread qui est concerné.
- ThreadIdx.xy représentant l'Id (selon x ou y) du thread qui est concerné
Le fait que l'on utilise le specifier global traduit le fait que la fonction est:
- Exécutée par le GPU.
- Appelée par le CPU.
On obtient ce résultat :
On note que la matrice issue de l'addition est la meme que celle calculée sur le CPU, cependant sont temps de calcul est fortement réduit.
Nous voulons maintenant effectuer la multiplication de deux matrices. Le calcul de chaque coefficient de la matrice de sortie est le suivant :
On effectue ce calcul sur le CPU avec cette fonction
On obtient ainsi la matice suivante :
Les valeurs sont bonnes il y a juste des arrondis.
On se propose maintenant de réaliser la même fonction mais en CUDA et exécutable par le GPU.
Chaque ligne est représentée par un Block et chaque colonne est représenté par un ID de thread au sein de ces blocks.
On obtient ainsi la matrice suivante :
On observe que les matrices issues du CPU et du GPU sont identiques.
Nous avons aussi codé une fonction permettant de créer une matrice de taille NxP remplie de coefficients tous égaux à une valeur fixée val. Cela nous permettra par la suite de faire des vérification rapides que nos autres fonctions marchent bien, notamment quand il s'agira de notre fonction de convolution.
De plus on crée une fonction permettant de faire une matrice de kernel carré de dimension dim identitaire dont la valeur centrale est fixée par la variable val.
Ces kernels sont très pratiques pour vérifier que nos convolutions marchent bien en un coup d'oeil.
On obtient donc les 6 noyaux de taille 5x5 suivants :
Nous voulons mettre en place une convolution en 2 dimensions de notre image 32x32x1 d'entrée issue de MNIST. Nous voulons réaliser cette convolution sur GPU de facons à diminuer au minimum notre temps de calcul.
Dans notre cas on souhaite faire la convolution de cette image par 6 kernels de taille 5x5, nous obtiendrons donc une sortie de taille 28x28x6. Layer 2- Convolution avec 6 noyaux de convolution de taille 5x5. La taille résultante est donc de 6x28x28.
Voici la fonction de convolution que nous avons utilisée pour la suite :
__global__ void cudaConv2D(float* M, float* kernel, float* Mout, int M_line, int M_col, int kernel_size, int nb_kernel){
int offset = (kernel_size-1)/2;
int out_line = M_line - offset;
int out_col = M_col - offset;
//Convolution d'une matrice par un kernel
int lig = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float conv = 0.0;
if (lig < out_line && col < out_col){
int temp = M_line * M_col;
for (int k_line = 0; k_line < kernel_size; k_line++) {
for (int k_col = 0; k_col < kernel_size; k_col++) {
for (int n_k = 0; n_k < nb_kernel; n_k++){
conv += M[(lig + k_line) * M_col + col + k_col + n_k * temp] * kernel[k_line * kernel_size + k_col + n_k * nb_kernel];
}
}
}
Mout[lig * out_col + col] = conv;
}
}Nous nous fixons maintenant l'objectif de faire un mean pooling 2x2 de la sortie de notre couche de convolution.
Cette étape permet de faire un subsampling de la feature map tout en introduisant une invariance à la translation. Ce type d'opération est très souvent utilisé dans après une couche de convolution.
Afin de vérifier notre fonction de convolution, nous avons choisi de faire la convolution d'une matrice de 32321 remplie de 1 par des noyaux identité dont la valeur centrale est 10.
Voici la matrice de sortie de la convolution :
Les résultats sont bien conformes à ce que nous attendions à avoir, à savoir une matrice 28286 remplie de 10.
Afin d'achever cette partie nous allons coder une fonction d'activation afin de l'appliquer en sortie de nos deux couches à chacun des coefficients de la matrice. Le choix s'est porté sur une fonction tanh :
Cette fonction renvoie une valeur entre -1 et 1 et celle-ci sature à 1 en à partir de 2 et à -1 à partir de -2; C'est pour cette raison que nous allons tester cette fonction à l'aide d'une matrice remplie de valeur entre -1 et 1 que nous allons convoluer avec un kenel unitaire suivi d'un average pooling.
On obtient la matrice suivante de taille 14x14 :
Celle-ci correspond bien à ce qu'on s'attend à avoir.
En ouvrant le LeNet5 on se rend compte que le résumé du modèle sous tensorFlow est le suivant :
Il nous manque une couche de pooling, une couche de conv ainsi que 3 couches Dense. nous disposons déjà des fonctions pour faire les couches conv et de pooling. Il faut donc que nous réalisions la fonction pour faire les couches dense et l'activation softmax. Nous avons donc écrit les fonctions suivantes :
- fonction dense :
__global__ void cudaFullyConnected(float *x_in, float *w, float *Mout, int lig_prec, int col_prec,int prof_prec){
int lig = threadIdx.x;
float temp = 0;
for (int i=0; i<lig_prec*col_prec*prof_prec;i++){
temp+=x_in[i]*w[i*lig_prec*col_prec*prof_prec +lig];
}
Mout[lig] = temp;
}N'arrivant pas à extraire les poids du modèle python de facons convenable, nous n'avons pas finalisé le modèle. Cependant nous avons quand mème réalisé la structure de celui-ci dans notre dernier fichier.