Add VGG16 architecture

src/cnn/backpropagation.c

@@ -3,11 +3,13 @@
 #include <math.h>
 
 #include "include/backpropagation.h"
+#include "../common/include/colors.h"
 #include "../common/include/utils.h"
 #include "include/struct.h"
 
 #include "include/config.h"
 
+
 /*
 * Softmax backward MSE
 */
@@ -101,7 +103,7 @@ void softmax_backward_cross_entropy(float* input, float* output, int size) {
 * Backward average pooling
 */
 #ifdef __CUDACC__
-__global__ void backward_average_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int n, int size) {
+__global__ void backward_average_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Équivalents respectifs de i, j et k dans la boucle effectuée par le cpu
     int idx = threadIdx.x + blockDim.x*blockIdx.x; // < depth
     int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width
@@ -110,45 +112,54 @@ __global__ void backward_average_pooling_kernel(float*** input, float*** output,
     if (idx >= depth || idy >= output_width || idz >= output_width) {
         return;
     }
+    int max_move = kernel_size - padding;
 
-    for (int a=0; a < size; a++) {
-        for (int b=0; b < size; b++) {
-            input[idx][size*idy +a][size*idz +b] += output[idx][idy][idz]/n;
+    for (int a=-padding; a < max_move; a++) {
+        for (int b=-padding; b < max_move; b++) {
+            int idy_2 = stride*idy +a;
+            int idz_2 = stride*idz +b;
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
+                int y = min(idy_2+1, min(kernel_size, input_width - idy_2));
+                int z = min(idz_2+1, min(kernel_size, input_width - idz_2));
+                input[idx][idy_2][idz_2] += output[idx][idy][idz]/(y*z);
+            }
         }
     }
 }
 
 
-void backward_average_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_average_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Make computation
     dim3 gridSize(i_div_up(depth, BLOCKSIZE_x), i_div_up(output_width, BLOCKSIZE_y), i_div_up(output_width, BLOCKSIZE_z));
     dim3 blockSize(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    int size = input_width/output_width; // Taille du pooling
-
     reset_3d_array(input, depth, input_width, input_width);
 
-    backward_average_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, size*size, size);
+    backward_average_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
 }
 #endif
 
-void backward_average_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_average_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     /* Input et output ont la même profondeur (depth) */
 
-    int size = input_width/output_width; // Taille du pooling
-    int n = size*size; // Nombre d'éléments dans le pooling
-
     reset_3d_array(input, depth, input_width, input_width);
+    int max_move = kernel_size - padding;
 
     for (int i=0; i < depth; i++) {
         for (int j=0; j < output_width; j++) {
             for (int k=0; k < output_width; k++) {
-                for (int a=0; a < size; a++) {
-                    for (int b=0; b < size; b++) {
-                        input[i][size*j +a][size*k +b] += output[i][j][k]/n;
+                for (int a=-padding; a < max_move; a++) {
+                    for (int b=-padding; b < max_move; b++) {
+                        int j_2 = stride*j +a;
+                        int k_2 = stride*k + b;
+                        if (not_outside(j_2, k_2, 0, input_width)){
+                            int j_3 = min(j_2+1, min(kernel_size, input_width - j_2));
+                            int k_3 = min(k_2+1, min(kernel_size, input_width - k_2));
+                            input[i][j_2][k_2] += output[i][j][k]/(j_3*k_3);
+                        }
                     }
                 }
             }
@@ -159,11 +170,11 @@ void backward_average_pooling_cpu(float*** input, float*** output, int input_wid
 #ifdef __CUDACC__
 extern "C"
 #endif
-void backward_average_pooling(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_average_pooling(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     #ifndef __CUDACC__
-    backward_average_pooling_cpu(input, output, input_width, output_width, depth);
+    backward_average_pooling_cpu(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     #else
-    backward_average_pooling_device(input, output, input_width, output_width, depth);
+    backward_average_pooling_device(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     #endif
 }
 
@@ -172,7 +183,7 @@ void backward_average_pooling(float*** input, float*** output, int input_width,
 * Backward max pooling
 */
 #ifdef __CUDACC__
-__global__ void backward_max_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int n, int size) {
+__global__ void backward_max_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Équivalents respectifs de i, j et k dans la boucle effectuée par le cpu
     int idx = threadIdx.x + blockDim.x*blockIdx.x; // < depth
     int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width
@@ -181,44 +192,51 @@ __global__ void backward_max_pooling_kernel(float*** input, float*** output, int
     if (idx >= depth || idy >= output_width || idz >= output_width) {
         return;
     }
-
+    int max_move = kernel_size - padding;
     float m = -FLT_MAX;
     int a_max = -1;
     int b_max = -1;
+    int cpt = 0;
 
-    for (int a=0; a < size; a++) {
-        for (int b=0; b < size; b++) {
-            if (input[idx][size*idy +a][size*idz +b] > m) {
-                m = input[idx][size*idy +a][size*idz +b];
-                a_max = a;
-                b_max = b;
+    for (int a=-padding; a < max_move; a++) {
+        for (int b=-padding; b < max_move; b++) {
+            int idy_2 = stride*idy +a;
+            int idz_2 = stride*idz +b;
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
+                if (input[idx][idy_2][idz_2] > m) {
+                    m = input[idx][idy_2][idz_2];
+                    a_max = a;
+                    b_max = b;
+                }
+                input[idx][idy_2][idz_2] = 0;
+                cpt++;
             }
-            input[idx][size*idy +a][size*idz +b] = 0;
         }
     }
-    input[idx][size*idy +a_max][size*idz +b_max] = output[idx][idy][idz]/n;
+    if (cpt==0) {
+        printf(RED "[ERROR]" RESET " Dimensions ou stride ou padding erroné dans 'backward_max_pooling_cpu'\n");
+    }
+    input[idx][stride*idy +a_max][stride*idz +b_max] = output[idx][idy][idz]/cpt;
 }
 
 
-void backward_max_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_max_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Make computation
     dim3 gridSize(i_div_up(depth, BLOCKSIZE_x), i_div_up(output_width, BLOCKSIZE_y), i_div_up(output_width, BLOCKSIZE_z));
     dim3 blockSize(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    int size = input_width/output_width; // Taille du pooling
-
-    backward_max_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, size*size, size);
+    backward_max_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
 }
 #endif
 
-void backward_max_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth) {
-    int size = input_width/output_width;
-
+void backward_max_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     float m; // Maximum
     int a_max, b_max; // Indices du maximum
+    int cpt;
+    int max_move = kernel_size - padding;
 
     for (int i=0; i < depth; i++) {
         for (int j=0; j < output_width; j++) {
@@ -226,18 +244,29 @@ void backward_max_pooling_cpu(float*** input, float*** output, int input_width,
                 m = -FLT_MAX;
                 a_max = -1;
                 b_max = -1;
+                cpt = 0;
 
-                for (int a=0; a < size; a++) {
-                    for (int b=0; b < size; b++) {
-                        if (input[i][size*j +a][size*k +b] > m) {
-                            m = input[i][size*j +a][size*k +b];
-                            a_max = a;
-                            b_max = b;
+                for (int a=-padding; a < max_move; a++) {
+                    for (int b=-padding; b < max_move; b++) {
+                        int j_2 = stride*j +a;
+                        int k_2 = stride*k +b;
+                        if (not_outside(j_2, k_2, 0, input_width)) {
+                            if (input[i][j_2][k_2] > m) {
+                                m = input[i][j_2][k_2];
+                                a_max = a;
+                                b_max = b;
+                            }
+                            input[i][j_2][k_2] = 0;
+                            cpt++;
                         }
-                        input[i][size*j +a][size*k +b] = 0;
                     }
                 }
-                input[i][size*j +a_max][size*k +b_max] = output[i][j][k]/(size*size);
+                if (cpt==0) {
+                    printf_error((char*)"Dimensions ou stride ou padding erroné dans 'backward_max_pooling_cpu'\n");
+                }
+                else {
+                    input[i][stride*j +a_max][stride*k +b_max] = output[i][j][k]/cpt;
+                }
             }
         }
     }
@@ -246,11 +275,11 @@ void backward_max_pooling_cpu(float*** input, float*** output, int input_width,
 #ifdef __CUDACC__
 extern "C"
 #endif
-void backward_max_pooling(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_max_pooling(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     #ifndef __CUDACC__
-    backward_max_pooling_cpu(input, output, input_width, output_width, depth);
+    backward_max_pooling_cpu(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     #else
-    backward_max_pooling_device(input, output, input_width, output_width, depth);
+    backward_max_pooling_device(input, output, input_width, output_width, kernel_size, depth, stride, padding);
     #endif
 }
 
@@ -480,15 +509,15 @@ __global__ void backward_convolution_dbias_kernel(Kernel_cnn* ker, float*** outp
     ker->d_bias[idx][idy][idz] += output[idx][idy][idz];
 }
 
-__global__ void backward_convolution_dweight_kernel(Kernel_cnn* ker, float*** input, float*** output, int input_depth, int output_depth, int output_width, int k_size) {
+__global__ void backward_convolution_dweight_kernel(Kernel_cnn* ker, float*** input, float*** output, int input_depth, int output_depth, int output_width, int kernel_size) {
     int idx = threadIdx.x + blockDim.x*blockIdx.x;
     int idy = threadIdx.y + blockDim.y*blockIdx.y;
     int idz = threadIdx.z + blockDim.z*blockIdx.z;
 
-    int idz1 = idz / k_size;
-    int idz2 = idz % k_size;
+    int idz1 = idz / kernel_size;
+    int idz2 = idz % kernel_size;
     
-    if (idx >= input_depth || idy >= output_depth || idz1 >= k_size || idz2 >= k_size) {
+    if (idx >= input_depth || idy >= output_depth || idz1 >= kernel_size || idz2 >= kernel_size) {
         return;
     }
     
@@ -526,23 +555,20 @@ __global__ void backward_convolution_propagate_kernel(Kernel_cnn* ker, float***
     input[idx][idy][idz] = tmp*( (*d_f)(input_z[idx][idy][idz]) );
 }
 
-void backward_convolution_device(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first) {
+void backward_convolution_device(Kernel_cnn* kernel, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first, int kernel_size, int padding, int stride) {
     // Bias Kernel
     dim3 gridSize1(i_div_up(output_depth, BLOCKSIZE_x), i_div_up(output_width, BLOCKSIZE_y), i_div_up(output_width, BLOCKSIZE_y));
     dim3 blockSize1(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    backward_convolution_dbias_kernel<<<gridSize1, blockSize1>>>(ker, output, output_depth, output_width);
+    backward_convolution_dbias_kernel<<<gridSize1, blockSize1>>>(kernel, output, output_depth, output_width);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
 
-    // Weights Kernel
-    int k_size = input_width - output_width +1;
-
-    dim3 gridSize2(i_div_up(input_depth, BLOCKSIZE_x), i_div_up(output_depth, BLOCKSIZE_y), i_div_up(k_size*k_size, BLOCKSIZE_y));
+    dim3 gridSize2(i_div_up(input_depth, BLOCKSIZE_x), i_div_up(output_depth, BLOCKSIZE_y), i_div_up(kernel_size*kernel_size, BLOCKSIZE_y));
     dim3 blockSize2(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    backward_convolution_dweight_kernel<<<gridSize2, blockSize2>>>(ker, input, output, input_depth, output_depth, output_width, k_size);
+    backward_convolution_dweight_kernel<<<gridSize2, blockSize2>>>(kernel, input, output, input_depth, output_depth, output_width, kernel_size);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
@@ -554,7 +580,7 @@ void backward_convolution_device(Kernel_cnn* ker, float*** input, float*** input
 
         funcPtr d_function = get_activation_function_cuda(activation);
 
-        backward_convolution_propagate_kernel<<<gridSize3, blockSize3>>>(ker, input, input_z, output, input_depth, input_width, output_depth, k_size, d_function);
+        backward_convolution_propagate_kernel<<<gridSize3, blockSize3>>>(kernel, input, input_z, output, input_depth, input_width, output_depth, kernel_size, d_function);
     
         gpuErrchk( cudaPeekAtLastError() );
         gpuErrchk( cudaDeviceSynchronize() );
@@ -563,9 +589,10 @@ void backward_convolution_device(Kernel_cnn* ker, float*** input, float*** input
 #endif
 
 
-void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first) {
+void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first, int kernel_size, int padding, int stride) {
     
     funcPtr d_function = get_activation_function(activation);
+    int max_move = kernel_size - padding;
 
     // Bias
     for (int i=0; i < output_depth; i++) {
@@ -577,16 +604,16 @@ void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z,
     }
 
     // Weights    
-    int k_size = input_width - output_width +1;
-    
     for (int h=0; h < input_depth; h++) {
         for (int i=0; i < output_depth; i++) {
-            for (int j=0; j < k_size; j++) {
-                for (int k=0; k < k_size; k++) {
+            for (int j=-padding; j < max_move; j++) {
+                for (int k=-padding; k < max_move; k++) {
                     float tmp = 0;
                     for (int l=0; l < output_width; l++) {
                         for (int m=0; m < output_width; m++) {
-                            tmp += input[h][l+j][m+k]*output[i][l][m];
+                            if (not_outside(l*stride+j, m*stride+k, 0, input_width)) {
+                                tmp += input[h][l*stride+j][m*stride+k]*output[i][l][m];
+                            }
                         }
                     }
                     ker->d_weights[h][i][j][k] += tmp;
@@ -595,26 +622,35 @@ void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z,
         }
     }
 
-    // Input
+    // Input TODO
     if (is_first==1) // Pas besoin de backpropager dans l'input
         return;
-    int min_m, max_m, min_n, max_n;
     for (int i=0; i < input_depth; i++) {
         for (int j=0; j < input_width; j++) {
             for (int k=0; k < input_width; k++) {
-                float tmp = 0;
-                for (int l=0; l < output_depth; l++) {
-                    min_m = max(0, k_size-1-j);
-                    max_m = min(k_size, input_width - j);
-                    min_n = max(0, k_size-1-k);
-                    max_n = min(k_size, input_width-k);
-                    for (int m=min_m; m < max_m; m++) {
-                        for (int n=min_n; n < max_n; n++) {
-                            tmp += output[l][j-k_size+m+1][k-k_size+n+1]*ker->weights[i][l][m][n];
+                input[i][j][k] = 0;
+            }
+        }
+    }
+    for (int h=0; h < input_depth; h++) {
+        for (int i=0; i < output_depth; i++) {
+            for (int j=-padding; j < max_move; j++) {
+                for (int k=-padding; k < max_move; k++) {
+                    for (int l=0; l < output_width; l++) {
+                        for (int m=0; m < output_width; m++) {
+                            if (not_outside(l*stride+j, m*stride+k, 0, input_width)) {
+                                input[h][l*stride+j][m*stride+k] += output[i][l][m]*ker->weights[h][i][j+padding][k+padding];
+                            }
                         }
                     }
                 }
-                input[i][j][k] = tmp*d_function(input_z[i][j][k]);
+            }
+        }
+    }
+    for (int i=0; i < input_depth; i++) {
+        for (int j=0; j < input_width; j++) {
+            for (int k=0; k < input_width; k++) {
+                input[i][j][k] = input[i][j][k]*d_function(input_z[i][j][k]);
             }
         }
     }
@@ -623,10 +659,10 @@ void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z,
 #ifdef __CUDACC__
 extern "C"
 #endif
-void backward_convolution(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first) {
+void backward_convolution(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first, int kernel_size, int padding, int stride) {
     #ifndef __CUDACC__
-    backward_convolution_cpu(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first);
+    backward_convolution_cpu(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first, kernel_size, padding, stride);
     #else
-    backward_convolution_device(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first);
+    backward_convolution_device(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first, kernel_size, padding, stride);
     #endif
 }

src/cnn/backpropagation.cu

@@ -3,11 +3,13 @@
 #include <math.h>
 
 #include "include/backpropagation.h"
+#include "../common/include/colors.h"
 #include "../common/include/utils.h"
 #include "include/struct.h"
 
 #include "include/config.h"
 
+
 /*
 * Softmax backward MSE
 */
@@ -101,7 +103,7 @@ void softmax_backward_cross_entropy(float* input, float* output, int size) {
 * Backward average pooling
 */
 #ifdef __CUDACC__
-__global__ void backward_average_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int n, int size) {
+__global__ void backward_average_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Équivalents respectifs de i, j et k dans la boucle effectuée par le cpu
     int idx = threadIdx.x + blockDim.x*blockIdx.x; // < depth
     int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width
@@ -110,45 +112,54 @@ __global__ void backward_average_pooling_kernel(float*** input, float*** output,
     if (idx >= depth || idy >= output_width || idz >= output_width) {
         return;
     }
+    int max_move = kernel_size - padding;
 
-    for (int a=0; a < size; a++) {
-        for (int b=0; b < size; b++) {
-            input[idx][size*idy +a][size*idz +b] += output[idx][idy][idz]/n;
+    for (int a=-padding; a < max_move; a++) {
+        for (int b=-padding; b < max_move; b++) {
+            int idy_2 = stride*idy +a;
+            int idz_2 = stride*idz +b;
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
+                int y = min(idy_2+1, min(kernel_size, input_width - idy_2));
+                int z = min(idz_2+1, min(kernel_size, input_width - idz_2));
+                input[idx][idy_2][idz_2] += output[idx][idy][idz]/(y*z);
+            }
         }
     }
 }
 
 
-void backward_average_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_average_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Make computation
     dim3 gridSize(i_div_up(depth, BLOCKSIZE_x), i_div_up(output_width, BLOCKSIZE_y), i_div_up(output_width, BLOCKSIZE_z));
     dim3 blockSize(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    int size = input_width/output_width; // Taille du pooling
-
     reset_3d_array(input, depth, input_width, input_width);
 
-    backward_average_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, size*size, size);
+    backward_average_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
 }
 #endif
 
-void backward_average_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_average_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     /* Input et output ont la même profondeur (depth) */
 
-    int size = input_width/output_width; // Taille du pooling
-    int n = size*size; // Nombre d'éléments dans le pooling
-
     reset_3d_array(input, depth, input_width, input_width);
+    int max_move = kernel_size - padding;
 
     for (int i=0; i < depth; i++) {
         for (int j=0; j < output_width; j++) {
             for (int k=0; k < output_width; k++) {
-                for (int a=0; a < size; a++) {
-                    for (int b=0; b < size; b++) {
-                        input[i][size*j +a][size*k +b] += output[i][j][k]/n;
+                for (int a=-padding; a < max_move; a++) {
+                    for (int b=-padding; b < max_move; b++) {
+                        int j_2 = stride*j +a;
+                        int k_2 = stride*k + b;
+                        if (not_outside(j_2, k_2, 0, input_width)){
+                            int j_3 = min(j_2+1, min(kernel_size, input_width - j_2));
+                            int k_3 = min(k_2+1, min(kernel_size, input_width - k_2));
+                            input[i][j_2][k_2] += output[i][j][k]/(j_3*k_3);
+                        }
                     }
                 }
             }
@@ -159,11 +170,11 @@ void backward_average_pooling_cpu(float*** input, float*** output, int input_wid
 #ifdef __CUDACC__
 extern "C"
 #endif
-void backward_average_pooling(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_average_pooling(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     #ifndef __CUDACC__
-    backward_average_pooling_cpu(input, output, input_width, output_width, depth);
+    backward_average_pooling_cpu(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     #else
-    backward_average_pooling_device(input, output, input_width, output_width, depth);
+    backward_average_pooling_device(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     #endif
 }
 
@@ -172,7 +183,7 @@ void backward_average_pooling(float*** input, float*** output, int input_width,
 * Backward max pooling
 */
 #ifdef __CUDACC__
-__global__ void backward_max_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int n, int size) {
+__global__ void backward_max_pooling_kernel(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Équivalents respectifs de i, j et k dans la boucle effectuée par le cpu
     int idx = threadIdx.x + blockDim.x*blockIdx.x; // < depth
     int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width
@@ -181,44 +192,51 @@ __global__ void backward_max_pooling_kernel(float*** input, float*** output, int
     if (idx >= depth || idy >= output_width || idz >= output_width) {
         return;
     }
-
+    int max_move = kernel_size - padding;
     float m = -FLT_MAX;
     int a_max = -1;
     int b_max = -1;
+    int cpt = 0;
 
-    for (int a=0; a < size; a++) {
-        for (int b=0; b < size; b++) {
-            if (input[idx][size*idy +a][size*idz +b] > m) {
-                m = input[idx][size*idy +a][size*idz +b];
-                a_max = a;
-                b_max = b;
+    for (int a=-padding; a < max_move; a++) {
+        for (int b=-padding; b < max_move; b++) {
+            int idy_2 = stride*idy +a;
+            int idz_2 = stride*idz +b;
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
+                if (input[idx][idy_2][idz_2] > m) {
+                    m = input[idx][idy_2][idz_2];
+                    a_max = a;
+                    b_max = b;
+                }
+                input[idx][idy_2][idz_2] = 0;
+                cpt++;
             }
-            input[idx][size*idy +a][size*idz +b] = 0;
         }
     }
-    input[idx][size*idy +a_max][size*idz +b_max] = output[idx][idy][idz]/n;
+    if (cpt==0) {
+        printf(RED "[ERROR]" RESET " Dimensions ou stride ou padding erroné dans 'backward_max_pooling_cpu'\n");
+    }
+    input[idx][stride*idy +a_max][stride*idz +b_max] = output[idx][idy][idz]/cpt;
 }
 
 
-void backward_max_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_max_pooling_device(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     // Make computation
     dim3 gridSize(i_div_up(depth, BLOCKSIZE_x), i_div_up(output_width, BLOCKSIZE_y), i_div_up(output_width, BLOCKSIZE_z));
     dim3 blockSize(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    int size = input_width/output_width; // Taille du pooling
-
-    backward_max_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, size*size, size);
+    backward_max_pooling_kernel<<<gridSize, blockSize>>>(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
 }
 #endif
 
-void backward_max_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth) {
-    int size = input_width/output_width;
-
+void backward_max_pooling_cpu(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     float m; // Maximum
     int a_max, b_max; // Indices du maximum
+    int cpt;
+    int max_move = kernel_size - padding;
 
     for (int i=0; i < depth; i++) {
         for (int j=0; j < output_width; j++) {
@@ -226,18 +244,29 @@ void backward_max_pooling_cpu(float*** input, float*** output, int input_width,
                 m = -FLT_MAX;
                 a_max = -1;
                 b_max = -1;
+                cpt = 0;
 
-                for (int a=0; a < size; a++) {
-                    for (int b=0; b < size; b++) {
-                        if (input[i][size*j +a][size*k +b] > m) {
-                            m = input[i][size*j +a][size*k +b];
-                            a_max = a;
-                            b_max = b;
+                for (int a=-padding; a < max_move; a++) {
+                    for (int b=-padding; b < max_move; b++) {
+                        int j_2 = stride*j +a;
+                        int k_2 = stride*k +b;
+                        if (not_outside(j_2, k_2, 0, input_width)) {
+                            if (input[i][j_2][k_2] > m) {
+                                m = input[i][j_2][k_2];
+                                a_max = a;
+                                b_max = b;
+                            }
+                            input[i][j_2][k_2] = 0;
+                            cpt++;
                         }
-                        input[i][size*j +a][size*k +b] = 0;
                     }
                 }
-                input[i][size*j +a_max][size*k +b_max] = output[i][j][k]/(size*size);
+                if (cpt==0) {
+                    printf_error((char*)"Dimensions ou stride ou padding erroné dans 'backward_max_pooling_cpu'\n");
+                }
+                else {
+                    input[i][stride*j +a_max][stride*k +b_max] = output[i][j][k]/cpt;
+                }
             }
         }
     }
@@ -246,11 +275,11 @@ void backward_max_pooling_cpu(float*** input, float*** output, int input_width,
 #ifdef __CUDACC__
 extern "C"
 #endif
-void backward_max_pooling(float*** input, float*** output, int input_width, int output_width, int depth) {
+void backward_max_pooling(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding) {
     #ifndef __CUDACC__
-    backward_max_pooling_cpu(input, output, input_width, output_width, depth);
+    backward_max_pooling_cpu(input, output, input_width, output_width, depth, kernel_size, stride, padding);
     #else
-    backward_max_pooling_device(input, output, input_width, output_width, depth);
+    backward_max_pooling_device(input, output, input_width, output_width, kernel_size, depth, stride, padding);
     #endif
 }
 
@@ -480,15 +509,15 @@ __global__ void backward_convolution_dbias_kernel(Kernel_cnn* ker, float*** outp
     ker->d_bias[idx][idy][idz] += output[idx][idy][idz];
 }
 
-__global__ void backward_convolution_dweight_kernel(Kernel_cnn* ker, float*** input, float*** output, int input_depth, int output_depth, int output_width, int k_size) {
+__global__ void backward_convolution_dweight_kernel(Kernel_cnn* ker, float*** input, float*** output, int input_depth, int output_depth, int output_width, int kernel_size) {
     int idx = threadIdx.x + blockDim.x*blockIdx.x;
     int idy = threadIdx.y + blockDim.y*blockIdx.y;
     int idz = threadIdx.z + blockDim.z*blockIdx.z;
 
-    int idz1 = idz / k_size;
-    int idz2 = idz % k_size;
+    int idz1 = idz / kernel_size;
+    int idz2 = idz % kernel_size;
     
-    if (idx >= input_depth || idy >= output_depth || idz1 >= k_size || idz2 >= k_size) {
+    if (idx >= input_depth || idy >= output_depth || idz1 >= kernel_size || idz2 >= kernel_size) {
         return;
     }
     
@@ -526,23 +555,20 @@ __global__ void backward_convolution_propagate_kernel(Kernel_cnn* ker, float***
     input[idx][idy][idz] = tmp*( (*d_f)(input_z[idx][idy][idz]) );
 }
 
-void backward_convolution_device(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first) {
+void backward_convolution_device(Kernel_cnn* kernel, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first, int kernel_size, int padding, int stride) {
     // Bias Kernel
     dim3 gridSize1(i_div_up(output_depth, BLOCKSIZE_x), i_div_up(output_width, BLOCKSIZE_y), i_div_up(output_width, BLOCKSIZE_y));
     dim3 blockSize1(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    backward_convolution_dbias_kernel<<<gridSize1, blockSize1>>>(ker, output, output_depth, output_width);
+    backward_convolution_dbias_kernel<<<gridSize1, blockSize1>>>(kernel, output, output_depth, output_width);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
 
-    // Weights Kernel
-    int k_size = input_width - output_width +1;
-
-    dim3 gridSize2(i_div_up(input_depth, BLOCKSIZE_x), i_div_up(output_depth, BLOCKSIZE_y), i_div_up(k_size*k_size, BLOCKSIZE_y));
+    dim3 gridSize2(i_div_up(input_depth, BLOCKSIZE_x), i_div_up(output_depth, BLOCKSIZE_y), i_div_up(kernel_size*kernel_size, BLOCKSIZE_y));
     dim3 blockSize2(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
 
-    backward_convolution_dweight_kernel<<<gridSize2, blockSize2>>>(ker, input, output, input_depth, output_depth, output_width, k_size);
+    backward_convolution_dweight_kernel<<<gridSize2, blockSize2>>>(kernel, input, output, input_depth, output_depth, output_width, kernel_size);
     
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
@@ -554,7 +580,7 @@ void backward_convolution_device(Kernel_cnn* ker, float*** input, float*** input
 
         funcPtr d_function = get_activation_function_cuda(activation);
 
-        backward_convolution_propagate_kernel<<<gridSize3, blockSize3>>>(ker, input, input_z, output, input_depth, input_width, output_depth, k_size, d_function);
+        backward_convolution_propagate_kernel<<<gridSize3, blockSize3>>>(kernel, input, input_z, output, input_depth, input_width, output_depth, kernel_size, d_function);
     
         gpuErrchk( cudaPeekAtLastError() );
         gpuErrchk( cudaDeviceSynchronize() );
@@ -563,9 +589,10 @@ void backward_convolution_device(Kernel_cnn* ker, float*** input, float*** input
 #endif
 
 
-void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first) {
+void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first, int kernel_size, int padding, int stride) {
     
     funcPtr d_function = get_activation_function(activation);
+    int max_move = kernel_size - padding;
 
     // Bias
     for (int i=0; i < output_depth; i++) {
@@ -577,16 +604,16 @@ void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z,
     }
 
     // Weights    
-    int k_size = input_width - output_width +1;
-    
     for (int h=0; h < input_depth; h++) {
         for (int i=0; i < output_depth; i++) {
-            for (int j=0; j < k_size; j++) {
-                for (int k=0; k < k_size; k++) {
+            for (int j=-padding; j < max_move; j++) {
+                for (int k=-padding; k < max_move; k++) {
                     float tmp = 0;
                     for (int l=0; l < output_width; l++) {
                         for (int m=0; m < output_width; m++) {
-                            tmp += input[h][l+j][m+k]*output[i][l][m];
+                            if (not_outside(l*stride+j, m*stride+k, 0, input_width)) {
+                                tmp += input[h][l*stride+j][m*stride+k]*output[i][l][m];
+                            }
                         }
                     }
                     ker->d_weights[h][i][j][k] += tmp;
@@ -595,26 +622,35 @@ void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z,
         }
     }
 
-    // Input
+    // Input TODO
     if (is_first==1) // Pas besoin de backpropager dans l'input
         return;
-    int min_m, max_m, min_n, max_n;
     for (int i=0; i < input_depth; i++) {
         for (int j=0; j < input_width; j++) {
             for (int k=0; k < input_width; k++) {
-                float tmp = 0;
-                for (int l=0; l < output_depth; l++) {
-                    min_m = max(0, k_size-1-j);
-                    max_m = min(k_size, input_width - j);
-                    min_n = max(0, k_size-1-k);
-                    max_n = min(k_size, input_width-k);
-                    for (int m=min_m; m < max_m; m++) {
-                        for (int n=min_n; n < max_n; n++) {
-                            tmp += output[l][j-k_size+m+1][k-k_size+n+1]*ker->weights[i][l][m][n];
+                input[i][j][k] = 0;
+            }
+        }
+    }
+    for (int h=0; h < input_depth; h++) {
+        for (int i=0; i < output_depth; i++) {
+            for (int j=-padding; j < max_move; j++) {
+                for (int k=-padding; k < max_move; k++) {
+                    for (int l=0; l < output_width; l++) {
+                        for (int m=0; m < output_width; m++) {
+                            if (not_outside(l*stride+j, m*stride+k, 0, input_width)) {
+                                input[h][l*stride+j][m*stride+k] += output[i][l][m]*ker->weights[h][i][j+padding][k+padding];
+                            }
                         }
                     }
                 }
-                input[i][j][k] = tmp*d_function(input_z[i][j][k]);
+            }
+        }
+    }
+    for (int i=0; i < input_depth; i++) {
+        for (int j=0; j < input_width; j++) {
+            for (int k=0; k < input_width; k++) {
+                input[i][j][k] = input[i][j][k]*d_function(input_z[i][j][k]);
             }
         }
     }
@@ -623,10 +659,10 @@ void backward_convolution_cpu(Kernel_cnn* ker, float*** input, float*** input_z,
 #ifdef __CUDACC__
 extern "C"
 #endif
-void backward_convolution(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first) {
+void backward_convolution(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first, int kernel_size, int padding, int stride) {
     #ifndef __CUDACC__
-    backward_convolution_cpu(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first);
+    backward_convolution_cpu(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first, kernel_size, padding, stride);
     #else
-    backward_convolution_device(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first);
+    backward_convolution_device(ker, input, input_z, output, input_depth, input_width, output_depth, output_width, activation, is_first, kernel_size, padding, stride);
     #endif
 }

src/cnn/cnn.c

@@ -258,10 +258,13 @@ void backward_propagation(Network* network, int wanted_number) {
 
         int is_last_layer = i==0;
         int activation = is_last_layer?SIGMOID:network->kernel[i-1]->activation;
+        int padding = k_i->padding;
+        int stride = k_i->stride;
 
 
         if (k_i->cnn) { // Convolution
-            backward_convolution(k_i->cnn, input, input_z, output, input_depth, input_width, output_depth, output_width, -activation, is_last_layer);
+            int kernel_size = k_i->cnn->k_size;
+            backward_convolution(k_i->cnn, input, input_z, output, input_depth, input_width, output_depth, output_width, -activation, is_last_layer, kernel_size, padding, stride);
         } else if (k_i->nn) { // Full connection
             if (k_i->linearisation == DOESNT_LINEARISE) { // Vecteur -> Vecteur
                 backward_dense(k_i->nn, input[0][0], input_z[0][0], output[0][0], input_width, output_width, -activation, is_last_layer);
@@ -269,10 +272,11 @@ void backward_propagation(Network* network, int wanted_number) {
                 backward_linearisation(k_i->nn, input, input_z, output[0][0], input_depth, input_width, output_width, -activation);
             }
         } else { // Pooling
+            int kernel_size = 2*padding + input_width + stride - output_width*stride;
             if (k_i->pooling == AVG_POOLING) {
-                backward_average_pooling(input, output, input_width, output_width, input_depth); // Depth pour input et output a la même valeur
+                backward_average_pooling(input, output, input_width, output_width, input_depth, kernel_size, stride, padding); // Depth pour input et output a la même valeur
             } else {
-                backward_max_pooling(input, output, input_width, output_width, input_depth); // Depth pour input et output a la même valeur
+                backward_max_pooling(input, output, input_width, output_width, input_depth, kernel_size, stride, padding); // Depth pour input et output a la même valeur
             }
         }
     }

src/cnn/convolution.c

@@ -7,12 +7,6 @@
 
 #include "include/config.h"
 
-#ifdef __CUDACC__
-__host__ __device__
-#endif
-int convolution_not_outside(int x, int y, int lower_bound, int upper_bound) {
-    return !(x < lower_bound || y < lower_bound || x >= upper_bound || y>= upper_bound);
-}
 
 void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, int output_width, int stride, int padding) {
     // c'est le kernel de input
@@ -34,7 +28,7 @@ void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, i
                         for (int c=-padding; c < max_move; c++) { // colonne du filtre
                             int x = (stride*j+b);
                             int y = (stride*k+c);
-                            if (convolution_not_outside(x, y, 0, input_width)) {
+                            if (not_outside(x, y, 0, input_width)) {
                                 f += kernel->weights[a][i][b][c]*input[a][stride*j+b][stride*k+c];
                             }
                         }
@@ -67,7 +61,7 @@ __global__ void make_convolution_kernel(Kernel_cnn* kernel, float*** input, floa
             for (int c=-padding; c < max_move; c++) {
                 int idy_2 = idy*stride+b;
                 int idz_2 = idz*stride+c;
-                if (convolution_not_outside(idy_2, idz_2, 0, input_width)) {
+                if (not_outside(idy_2, idz_2, 0, input_width)) {
                     f += kernel->weights[a][idx][b][c]*input[a][idy_2][idz_2];
                 }
             }

src/cnn/convolution.cu

@@ -7,16 +7,6 @@
 
 #include "include/config.h"
 
-#ifdef __CUDACC__
-__host__ __device__
-#endif
-int convolution_not_outside(int x, int y, int lower_bound, int upper_bound) {
-    // On renvoie true si et seulement si _ et _:
-    // lower_bound <= x < upper_bound
-    // lower_bound <= y < upper_bound
-    
-    return !(x < lower_bound || y < lower_bound || x >= upper_bound || y>= upper_bound);
-}
 
 void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, int output_width, int stride, int padding) {
     // c'est le kernel de input
@@ -38,7 +28,7 @@ void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, i
                         for (int c=-padding; c < max_move; c++) { // colonne du filtre
                             int x = (stride*j+b);
                             int y = (stride*k+c);
-                            if (convolution_not_outside(x, y, 0, input_width)) {
+                            if (not_outside(x, y, 0, input_width)) {
                                 f += kernel->weights[a][i][b][c]*input[a][stride*j+b][stride*k+c];
                             }
                         }
@@ -71,7 +61,7 @@ __global__ void make_convolution_kernel(Kernel_cnn* kernel, float*** input, floa
             for (int c=-padding; c < max_move; c++) {
                 int idy_2 = idy*stride+b;
                 int idz_2 = idz*stride+c;
-                if (convolution_not_outside(idy_2, idz_2, 0, input_width)) {
+                if (not_outside(idy_2, idz_2, 0, input_width)) {
                     f += kernel->weights[a][idx][b][c]*input[a][idy_2][idz_2];
                 }
             }

src/cnn/creation.c

@@ -9,7 +9,7 @@
 
 #include "include/creation.h"
 
-Network* create_network(int max_size, float learning_rate, int dropout, int activation, int initialisation, int input_width, int input_depth) {
+Network* create_network(int max_size, float learning_rate, int dropout, int initialisation, int input_width, int input_depth) {
     if (dropout < 0 || dropout > 100) {
         printf_error("La probabilité de dropout n'est pas respecté, elle doit être comprise entre 0 et 100\n");
     }
@@ -27,19 +27,15 @@ Network* create_network(int max_size, float learning_rate, int dropout, int acti
     for (int i=0; i < max_size-1; i++) {
         network->kernel[i] = (Kernel*)nalloc(1, sizeof(Kernel));
     }
-    network->kernel[0]->linearisation = DOESNT_LINEARISE;
-    network->kernel[0]->activation = activation;
     network->width[0] = input_width;
     network->depth[0] = input_depth;
-    network->kernel[0]->nn = NULL;
-    network->kernel[0]->cnn = NULL;
     create_a_cube_input_layer(network, 0, input_depth, input_width);
     create_a_cube_input_z_layer(network, 0, input_depth, input_width);
     return network;
 }
 
 Network* create_network_lenet5(float learning_rate, int dropout, int activation, int initialisation, int input_width, int input_depth) {
-    Network* network = create_network(8, learning_rate, dropout, activation, initialisation, input_width, input_depth);
+    Network* network = create_network(8, learning_rate, dropout, initialisation, input_width, input_depth);
     add_convolution(network, 5, 6, 1, 0, activation);
     add_average_pooling(network, 2, 2, 0);
     add_convolution(network, 5, 16, 1, 0, activation);
@@ -50,8 +46,56 @@ Network* create_network_lenet5(float learning_rate, int dropout, int activation,
     return network;
 }
 
+Network* create_network_alexnet(float learning_rate, int dropout, int activation, int initialisation, int size_output) {
+    Network* network = create_network(12, learning_rate, dropout, initialisation, 227, 3);
+    add_convolution(network, 11, 96, 4, 0, activation);
+    add_average_pooling(network, 3, 2, 0);
+    add_convolution(network, 5, 256, 1, 2, activation);
+    add_average_pooling(network, 3, 2, 0);
+    add_convolution(network, 3, 384, 1, 1, activation);
+    add_convolution(network, 3, 384, 1, 1, activation);
+    add_convolution(network, 3, 256, 1, 1, activation);
+    add_average_pooling(network, 3, 2, 0);
+    add_dense_linearisation(network, 4096, activation);
+    add_dense(network, 4096, activation);
+    add_dense(network, size_output, SOFTMAX);
+    return network;
+}
+
+Network* create_network_VGG16(float learning_rate, int dropout, int activation, int initialisation, int size_output) {
+    Network* network = create_network(23, learning_rate, dropout, initialisation, 256, 3);
+    add_convolution(network, 3, 64, 1, 0, activation); // Conv3-64
+    add_convolution(network, 3, 64, 1, 0, activation); // Conv3-64
+    add_average_pooling(network, 2, 2, 0); // Max Pool
+
+    add_convolution(network, 3, 128, 1, 0, activation); // Conv3-128
+    add_convolution(network, 1, 128, 1, 0, activation); // Conv1-128
+    add_average_pooling(network, 2, 2, 0); // Max Pool
+
+    add_convolution(network, 3, 256, 1, 0, activation); // Conv3-256
+    add_convolution(network, 3, 256, 1, 0, activation); // Conv3-256
+    add_convolution(network, 1, 256, 1, 0, activation); // Conv1-256
+    add_average_pooling(network, 2, 2, 0); // Max Pool
+
+    add_convolution(network, 3, 512, 1, 0, activation); // Conv3-512
+    add_convolution(network, 3, 512, 1, 0, activation); // Conv3-512
+    add_convolution(network, 1, 512, 1, 0, activation); // Conv1-512
+    add_average_pooling(network, 2, 2, 0); // Max Pool
+
+    add_convolution(network, 3, 512, 1, 0, activation); // Conv3-512
+    add_convolution(network, 3, 512, 1, 0, activation); // Conv3-512
+    add_convolution(network, 1, 512, 1, 0, activation); // Conv1-512
+    add_average_pooling(network, 2, 2, 0); // Max Pool
+
+    add_dense_linearisation(network, 2048, activation);
+    add_dense(network, 2048, activation);
+    add_dense(network, 256, activation);
+    add_dense(network, size_output, SOFTMAX);
+    return network;
+}
+
 Network* create_simple_one(float learning_rate, int dropout, int activation, int initialisation, int input_width, int input_depth) {
-    Network* network = create_network(3, learning_rate, dropout, activation, initialisation, input_width, input_depth);
+    Network* network = create_network(3, learning_rate, dropout, initialisation, input_width, input_depth);
     add_dense_linearisation(network, 80, activation);
     add_dense(network, 10, SOFTMAX);
     return network;
@@ -132,6 +176,8 @@ void add_max_pooling(Network* network, int kernel_size, int stride, int padding)
 
     network->kernel[k_pos]->cnn = NULL;
     network->kernel[k_pos]->nn = NULL;
+    network->kernel[k_pos]->stride = stride;
+    network->kernel[k_pos]->padding = padding;
     network->kernel[k_pos]->activation = IDENTITY; // Ne contient pas de fonction d'activation
     network->kernel[k_pos]->linearisation = DOESNT_LINEARISE;
     network->kernel[k_pos]->pooling = MAX_POOLING;

src/cnn/include/backpropagation.h

@@ -31,7 +31,7 @@ extern "C"
 * Transfert les informations d'erreur à travers une couche d'average pooling
 * en considérant cross_entropy comme fonction d'erreur
 */
-void backward_average_pooling(float*** input, float*** output, int input_width, int output_width, int depth);
+void backward_average_pooling(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding);
 
 
 #ifdef __CUDACC__
@@ -41,7 +41,7 @@ extern "C"
 * Transfert les informations d'erreur à travers une couche de max pooling
 * en considérant cross_entropy comme fonction d'erreur
 */
-void backward_max_pooling(float*** input, float*** output, int input_width, int output_width, int depth);
+void backward_max_pooling(float*** input, float*** output, int input_width, int output_width, int depth, int kernel_size, int stride, int padding);
 
 
 #ifdef __CUDACC__
@@ -68,6 +68,6 @@ extern "C"
 /*
 * Transfert les informations d'erreur à travers un couche de convolution
 */
-void backward_convolution(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first);
+void backward_convolution(Kernel_cnn* ker, float*** input, float*** input_z, float*** output, int input_depth, int input_width, int output_depth, int output_width, int activation, int is_first, int kernel_size, int padding, int stride);
 
 #endif

src/cnn/include/convolution.h

@@ -1,14 +1,5 @@
 #include "struct.h"
 
-#ifdef __CUDACC__
-__host__ __device__
-#endif
-/*
-On renvoie true si et seulement si _ et _:
-lower_bound <= x < upper_bound
-lower_bound <= y < upper_bound
-*/
-int convolution_not_outside(int x, int y, int lower_bound, int upper_bound);
 
 /*
 * Effectue la convolution naïvement sur le processeur

src/cnn/include/creation.h

@@ -7,13 +7,25 @@
 /*
 * Créé un réseau qui peut contenir max_size couche (dont celle d'input et d'output)
 */
-Network* create_network(int max_size, float learning_rate, int dropout, int activation, int initialisation, int input_width, int input_depth);
+Network* create_network(int max_size, float learning_rate, int dropout, int initialisation, int input_width, int input_depth);
 
 /*
 * Renvoie un réseau suivant l'architecture LeNet5
 */
 Network* create_network_lenet5(float learning_rate, int dropout, int activation, int initialisation, int input_width, int input_depth);
 
+/*
+* Renvoie un réseau suivant l'architecture AlexNet
+* C'est à dire en entrée 3x227x227 et une sortie de taille 'size_output'
+*/
+Network* create_network_alexnet(float learning_rate, int dropout, int activation, int initialisation, int size_output);
+
+/*
+* Renvoie un réseau suivant l'architecture VGG16 modifiée pour prendre en entrée 3x256x256
+* et une sortie de taille 'size_output'
+*/
+Network* create_network_VGG16(float learning_rate, int dropout, int activation, int initialisation, int size_output);
+
 /*
 * Renvoie un réseau sans convolution, similaire à celui utilisé dans src/dense
 */

src/cnn/include/make.h

@@ -3,15 +3,6 @@
 #ifndef DEF_MAKE_H
 #define DEF_MAKE_H
 
-#ifdef __CUDACC__
-__host__ __device__
-#endif
-/*
-* On renvoie true si et seulement si _ et _:
-* lower_bound <= y < upper_bound
-* lower_bound <= x < upper_bound
-*/
-int pooling_not_outside(int x, int y, int lower_bound, int upper_bound);
 
 /*
 * Effectue la propagation d'une convolution avec stride et padding choisis sur le processeur

src/cnn/include/struct.h

@@ -12,18 +12,18 @@
 
 typedef struct Kernel_cnn {
     // Noyau ayant une couche matricielle en sortie
-    int k_size; // k_size = input_width - output_width + 1
+    int k_size; // k_size = 2*padding + input_width + stride - output_width*stride
     int rows; // Depth de l'input
     int columns; // Depth de l'output
 
-    float*** bias; // bias[columns][output_width][output_width] <=> bias[depth output][dim output][dim output]
+    float*** bias; // bias[columns][output_width][output_width] <=> bias[output depth][output width][output width]
     float*** d_bias; // d_bias[columns][output_width][output_width]
     #ifdef ADAM_CNN_BIAS
         float*** s_d_bias; // s_d_bias[columns][output_width][output_width]
         float*** v_d_bias; // v_d_bias[columns][output_width][output_width]
     #endif
 
-    float**** weights; // weights[rows][columns][k_size][k_size] <=> weights[depth input][depth output][size kernel][size kernel]
+    float**** weights; // weights[rows][columns][k_size][k_size] <=> weights[input depth][output depth][kernel size][kernel size]
     float**** d_weights; // d_weights[rows][columns][k_size][k_size]
     #ifdef ADAM_CNN_WEIGHTS
         float**** s_d_weights; // s_d_weights[rows][columns][k_size][k_size]

src/cnn/make.c

@@ -10,12 +10,6 @@
 
 #include "include/config.h"
 
-#ifdef __CUDACC__
-__host__ __device__
-#endif
-int pooling_not_outside(int x, int y, int lower_bound, int upper_bound) {
-    return !(x < lower_bound || y < lower_bound || x >= upper_bound || y>= upper_bound);
-}
 
 /* 
 * Average Pooling
@@ -40,7 +34,7 @@ __global__ void make_average_pooling_kernel(float*** input, float*** output, int
         for (int b=-padding; b < max_move; b++) {
             int idy_2 = stride*idy +a;
             int idz_2 = stride*idz +b;
-            if (pooling_not_outside(idy_2, idz_2, 0, input_width)) {
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
                 sum += input[idx][idy_2][idz_2];
                 nb_elements++;
             }
@@ -75,7 +69,7 @@ void make_average_pooling_cpu(float*** input, float*** output, int size, int out
                     for (int b=-padding; b < max_move; b++) {
                         int j_2 = stride*j +a;
                         int k_2 = stride*k +b;
-                        if (pooling_not_outside(j_2, k_2, 0, input_width)) {
+                        if (not_outside(j_2, k_2, 0, input_width)) {
                             sum += input[i][j_2][k_2];
                             nb_elements++;
                         }
@@ -125,7 +119,7 @@ __global__ void make_max_pooling_kernel(float*** input, float*** output, int siz
         for (int b=-padding; b < max_move; b++) {
             int idy_2 = stride*idy +a;
             int idz_2 = stride*idz +b;
-            if (pooling_not_outside(idy_2, idz_2, 0, input_width)) {
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
                 temp = input[idx][idy_2][idz_2];
                 m = m > temp ? m : temp; // max(m, temp)
             }
@@ -159,7 +153,7 @@ void make_max_pooling_cpu(float*** input, float*** output, int size, int output_
                     for (int b=-padding; b < max_move; b++) {
                         int j_2 = stride*j +a;
                         int k_2 = stride*k +b;
-                        if (pooling_not_outside(j_2, k_2, 0, input_width)) {
+                        if (not_outside(j_2, k_2, 0, input_width)) {
                             m = fmaxf(m, input[i][j_2][k_2]);
                         }
                     }

src/cnn/make.cu

@@ -10,12 +10,6 @@
 
 #include "include/config.h"
 
-#ifdef __CUDACC__
-__host__ __device__
-#endif
-int pooling_not_outside(int x, int y, int lower_bound, int upper_bound) {
-    return !(x < lower_bound || y < lower_bound || x >= upper_bound || y>= upper_bound);
-}
 
 /* 
 * Average Pooling
@@ -40,7 +34,7 @@ __global__ void make_average_pooling_kernel(float*** input, float*** output, int
         for (int b=-padding; b < max_move; b++) {
             int idy_2 = stride*idy +a;
             int idz_2 = stride*idz +b;
-            if (pooling_not_outside(idy_2, idz_2, 0, input_width)) {
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
                 sum += input[idx][idy_2][idz_2];
                 nb_elements++;
             }
@@ -75,7 +69,7 @@ void make_average_pooling_cpu(float*** input, float*** output, int size, int out
                     for (int b=-padding; b < max_move; b++) {
                         int j_2 = stride*j +a;
                         int k_2 = stride*k +b;
-                        if (pooling_not_outside(j_2, k_2, 0, input_width)) {
+                        if (not_outside(j_2, k_2, 0, input_width)) {
                             sum += input[i][j_2][k_2];
                             nb_elements++;
                         }
@@ -125,7 +119,7 @@ __global__ void make_max_pooling_kernel(float*** input, float*** output, int siz
         for (int b=-padding; b < max_move; b++) {
             int idy_2 = stride*idy +a;
             int idz_2 = stride*idz +b;
-            if (pooling_not_outside(idy_2, idz_2, 0, input_width)) {
+            if (not_outside(idy_2, idz_2, 0, input_width)) {
                 temp = input[idx][idy_2][idz_2];
                 m = m > temp ? m : temp; // max(m, temp)
             }
@@ -159,7 +153,7 @@ void make_max_pooling_cpu(float*** input, float*** output, int size, int output_
                     for (int b=-padding; b < max_move; b++) {
                         int j_2 = stride*j +a;
                         int k_2 = stride*k +b;
-                        if (pooling_not_outside(j_2, k_2, 0, input_width)) {
+                        if (not_outside(j_2, k_2, 0, input_width)) {
                             m = fmaxf(m, input[i][j_2][k_2]);
                         }
                     }

src/common/include/utils.h

@@ -1,5 +1,6 @@
-#include <stdio.h>
 #include <stdbool.h>
+#include <stdlib.h>
+#include <stdio.h>
 
 #ifdef USE_CUDA
    #ifndef __CUDACC__
@@ -39,6 +40,16 @@ int max(int a, int b);
 #endif
 
 
+#ifdef __CUDACC__
+__host__ __device__
+#endif
+/*
+* On renvoie true si et seulement si _ et _:
+* lower_bound <= x < upper_bound
+* lower_bound <= y < upper_bound
+*/
+bool not_outside(int x, int y, int lower_bound, int upper_bound);
+
 /*
 * Partie entière supérieure de a/b
 */

src/common/utils.c

@@ -1,3 +1,4 @@
+#include <stdbool.h>
 #include <stdlib.h>
 #include <stdio.h>
 #ifdef USE_CUDA
@@ -27,6 +28,13 @@ int max(int a, int b) {
 #endif
 
 
+#ifdef __CUDACC__
+__host__ __device__
+#endif
+bool not_outside(int x, int y, int lower_bound, int upper_bound) {
+    return !(x < lower_bound || y < lower_bound || x >= upper_bound || y>= upper_bound);
+}
+
 int i_div_up(int a, int b) { // Partie entière supérieure de a/b
     return ((a % b) != 0) ? (a / b + 1) : (a / b);
 }

src/common/utils.cu

@@ -1,3 +1,4 @@
+#include <stdbool.h>
 #include <stdlib.h>
 #include <stdio.h>
 #ifdef USE_CUDA
@@ -27,6 +28,13 @@ int max(int a, int b) {
 #endif
 
 
+#ifdef __CUDACC__
+__host__ __device__
+#endif
+bool not_outside(int x, int y, int lower_bound, int upper_bound) {
+    return !(x < lower_bound || y < lower_bound || x >= upper_bound || y>= upper_bound);
+}
+
 int i_div_up(int a, int b) { // Partie entière supérieure de a/b
     return ((a % b) != 0) ? (a / b + 1) : (a / b);
 }

test/cnn_structure.c

@@ -24,24 +24,27 @@ int main() {
             } else {
                 printf("\n==== Couche %d de type "YELLOW"Max Pooling"RESET" ====\n", i);
             }
+            int kernel_size = 2*kernel->padding + network->width[i] + kernel->stride - network->width[i+1]*kernel->stride;
+            printf("kernel: %dx%d, pad=%d, stride=%d\n", kernel_size, kernel_size, kernel->padding, kernel->stride);
         } else if (!kernel->cnn) {
             printf("\n==== Couche %d de type "GREEN"NN"RESET" ====\n", i);
+            if (kernel->linearisation) {
+                printf(YELLOW"Linéarisation: %d\n"RESET, kernel->linearisation);
+            }
             printf("input: %d\n", kernel->nn->size_input);
             printf("output: %d\n", kernel->nn->size_output);
         } else {
             printf("\n==== Couche %d de type "BLUE"CNN"RESET" ====\n", i);
-            printf("k_size: %d\n", kernel->cnn->k_size);
-            printf("rows: %d\n", kernel->cnn->rows);
-            printf("columns: %d\n", kernel->cnn->columns);
+            printf("kernel: %dx%d, pad=%d, stride=%d\n", kernel->cnn->k_size, kernel->cnn->k_size, kernel->padding, kernel->stride);
+            printf("%d kernels\n", kernel->cnn->columns);
         }
-        if (kernel->linearisation) {
-            printf(YELLOW"Linéarisation: %d\n"RESET, kernel->linearisation);
+        if (!kernel->nn) {
+            printf("depth: %d\n", network->depth[i]);
+            printf("width: %d\n", network->width[i]);
+        }
+        if (kernel->nn || kernel->cnn) {
+            printf("activation: %d\n", kernel->activation);
         }
-        printf("width: %d\n", network->width[i]);
-        printf("depth: %d\n", network->depth[i]);
-        printf("activation: %d\n", kernel->activation);
-        printf("stride: %d\n", kernel->stride);
-        printf("padding: %d\n", kernel->padding);
     }
     printf("\n" GREEN "OK\n" RESET);