From 3dd2e33fa9f91c49a62436ec4812569bf5c1942c Mon Sep 17 00:00:00 2001 From: julienChemillier Date: Sat, 13 May 2023 15:39:22 +0200 Subject: [PATCH] Fix issues with the network not converging --- src/cnn/make.c | 13 +++++++++---- src/cnn/make.cu | 13 +++++++++---- src/cnn/neuron_io.c | 26 ++++++++++++++++++++------ src/cnn/utils.c | 8 ++++++++ test/cnn_structure.c | 2 ++ 5 files changed, 48 insertions(+), 14 deletions(-) diff --git a/src/cnn/make.c b/src/cnn/make.c index 3a6572b..423eb99 100644 --- a/src/cnn/make.c +++ b/src/cnn/make.c @@ -24,6 +24,7 @@ __global__ void make_average_pooling_kernel(float*** input, float*** output, int int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width int idz = threadIdx.z + blockDim.z*blockIdx.z; // < output_width int max_move = size - padding; + int input_dim = output_width*stride - 2*padding + size - stride; if (idx >= output_depth || idy >= output_width || idz >= output_width) { return; @@ -36,8 +37,9 @@ __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, output_width)) { + if (pooling_not_outside(idy_2, idz_2, 0, input_dim)) { sum += input[idx][idy_2][idz_2]; + nb_elements++; } } } @@ -59,6 +61,7 @@ void make_average_pooling_cpu(float*** input, float*** output, int size, int out // input[output_depth][output_width+size-1][output_width+size-1] // output[output_depth][output_width][output_width] int max_move = size - padding; + int input_dim = output_width*stride - 2*padding + size - stride; for (int i=0; i < output_depth; i++) { for (int j=0; j < output_width; j++) { @@ -69,7 +72,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, output_width)) { + if (pooling_not_outside(j_2, k_2, 0, input_dim)) { sum += input[i][j_2][k_2]; nb_elements++; } @@ -105,6 +108,7 @@ __global__ void make_max_pooling_kernel(float*** input, float*** output, int siz int idx = threadIdx.x + blockDim.x*blockIdx.x; // < output_depth int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width int idz = threadIdx.z + blockDim.z*blockIdx.z; // < output_width + int input_dim = output_width*stride - 2*padding + size - stride; if (idx >= output_depth || idy >= output_width || idz >= output_width) { return; @@ -118,7 +122,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, output_width)) { + if (pooling_not_outside(idy_2, idz_2, 0, input_dim)) { temp = input[idx][idy_2][idz_2]; m = m > temp ? m : temp; // max(m, temp) } @@ -142,6 +146,7 @@ void make_max_pooling_cpu(float*** input, float*** output, int size, int output_ // input[output_depth][output_width+size-1][output_width+size-1] // output[output_depth][output_width][output_width] int max_move = size - padding; + int input_dim = output_width*stride - 2*padding + size - stride; float m; for (int i=0; i < output_depth; i++) { for (int j=0; j < output_width; j++) { @@ -151,7 +156,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, output_width)) { + if (pooling_not_outside(j_2, k_2, 0, input_dim)) { m = fmaxf(m, input[i][j_2][k_2]); } } diff --git a/src/cnn/make.cu b/src/cnn/make.cu index 3a6572b..423eb99 100644 --- a/src/cnn/make.cu +++ b/src/cnn/make.cu @@ -24,6 +24,7 @@ __global__ void make_average_pooling_kernel(float*** input, float*** output, int int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width int idz = threadIdx.z + blockDim.z*blockIdx.z; // < output_width int max_move = size - padding; + int input_dim = output_width*stride - 2*padding + size - stride; if (idx >= output_depth || idy >= output_width || idz >= output_width) { return; @@ -36,8 +37,9 @@ __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, output_width)) { + if (pooling_not_outside(idy_2, idz_2, 0, input_dim)) { sum += input[idx][idy_2][idz_2]; + nb_elements++; } } } @@ -59,6 +61,7 @@ void make_average_pooling_cpu(float*** input, float*** output, int size, int out // input[output_depth][output_width+size-1][output_width+size-1] // output[output_depth][output_width][output_width] int max_move = size - padding; + int input_dim = output_width*stride - 2*padding + size - stride; for (int i=0; i < output_depth; i++) { for (int j=0; j < output_width; j++) { @@ -69,7 +72,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, output_width)) { + if (pooling_not_outside(j_2, k_2, 0, input_dim)) { sum += input[i][j_2][k_2]; nb_elements++; } @@ -105,6 +108,7 @@ __global__ void make_max_pooling_kernel(float*** input, float*** output, int siz int idx = threadIdx.x + blockDim.x*blockIdx.x; // < output_depth int idy = threadIdx.y + blockDim.y*blockIdx.y; // < output_width int idz = threadIdx.z + blockDim.z*blockIdx.z; // < output_width + int input_dim = output_width*stride - 2*padding + size - stride; if (idx >= output_depth || idy >= output_width || idz >= output_width) { return; @@ -118,7 +122,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, output_width)) { + if (pooling_not_outside(idy_2, idz_2, 0, input_dim)) { temp = input[idx][idy_2][idz_2]; m = m > temp ? m : temp; // max(m, temp) } @@ -142,6 +146,7 @@ void make_max_pooling_cpu(float*** input, float*** output, int size, int output_ // input[output_depth][output_width+size-1][output_width+size-1] // output[output_depth][output_width][output_width] int max_move = size - padding; + int input_dim = output_width*stride - 2*padding + size - stride; float m; for (int i=0; i < output_depth; i++) { for (int j=0; j < output_width; j++) { @@ -151,7 +156,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, output_width)) { + if (pooling_not_outside(j_2, k_2, 0, input_dim)) { m = fmaxf(m, input[i][j_2][k_2]); } } diff --git a/src/cnn/neuron_io.c b/src/cnn/neuron_io.c index 44b7309..06c935e 100644 --- a/src/cnn/neuron_io.c +++ b/src/cnn/neuron_io.c @@ -76,12 +76,14 @@ void write_couche(Network* network, int indice_couche, int type_couche, FILE* pt int output_dim = network->width[indice_couche+1]; // Écriture du pré-corps - uint32_t pre_buffer[5]; + uint32_t pre_buffer[7]; pre_buffer[0] = kernel->activation; pre_buffer[1] = kernel->linearisation; pre_buffer[2] = cnn->k_size; pre_buffer[3] = cnn->rows; pre_buffer[4] = cnn->columns; + pre_buffer[5] = kernel->stride; + pre_buffer[6] = kernel->padding; fwrite(pre_buffer, sizeof(pre_buffer), 1, ptr); // Écriture du corps @@ -112,11 +114,13 @@ void write_couche(Network* network, int indice_couche, int type_couche, FILE* pt Kernel_nn* nn = kernel->nn; // Écriture du pré-corps - uint32_t pre_buffer[4]; + uint32_t pre_buffer[6]; pre_buffer[0] = kernel->activation; pre_buffer[1] = kernel->linearisation; pre_buffer[2] = nn->size_input; pre_buffer[3] = nn->size_output; + pre_buffer[4] = kernel->stride; + pre_buffer[5] = kernel->padding; fwrite(pre_buffer, sizeof(pre_buffer), 1, ptr); // Écriture du corps @@ -135,9 +139,11 @@ void write_couche(Network* network, int indice_couche, int type_couche, FILE* pt fwrite(buffer, sizeof(buffer), 1, ptr); } } else if (type_couche == 2) { // Cas du Pooling Layer - uint32_t pre_buffer[2]; + uint32_t pre_buffer[4]; pre_buffer[0] = kernel->linearisation; pre_buffer[1] = kernel->pooling; + pre_buffer[2] = kernel->stride; + pre_buffer[3] = kernel->padding; fwrite(pre_buffer, sizeof(pre_buffer), 1, ptr); } } @@ -234,7 +240,7 @@ Kernel* read_kernel(int type_couche, int output_dim, FILE* ptr) { // Lecture du "Pré-corps" kernel->cnn = (Kernel_cnn*)nalloc(1, sizeof(Kernel_cnn)); kernel->nn = NULL; - uint32_t buffer[5]; + uint32_t buffer[7]; (void) !fread(&buffer, sizeof(buffer), 1, ptr); kernel->activation = buffer[0]; @@ -242,6 +248,8 @@ Kernel* read_kernel(int type_couche, int output_dim, FILE* ptr) { kernel->cnn->k_size = buffer[2]; kernel->cnn->rows = buffer[3]; kernel->cnn->columns = buffer[4]; + kernel->stride = buffer[5]; + kernel->padding = buffer[6]; // Lecture du corps Kernel_cnn* cnn = kernel->cnn; @@ -322,13 +330,15 @@ Kernel* read_kernel(int type_couche, int output_dim, FILE* ptr) { // Lecture du "Pré-corps" kernel->nn = (Kernel_nn*)nalloc(1, sizeof(Kernel_nn)); kernel->cnn = NULL; - uint32_t buffer[4]; + uint32_t buffer[6]; (void) !fread(&buffer, sizeof(buffer), 1, ptr); kernel->activation = buffer[0]; kernel->linearisation = buffer[1]; kernel->nn->size_input = buffer[2]; kernel->nn->size_output = buffer[3]; + kernel->stride = buffer[4]; + kernel->padding = buffer[5]; // Lecture du corps Kernel_nn* nn = kernel->nn; @@ -374,15 +384,19 @@ Kernel* read_kernel(int type_couche, int output_dim, FILE* ptr) { } } } else if (type_couche == POOLING) { // Cas du Pooling Layer - uint32_t pooling, linearisation; + uint32_t pooling, linearisation, stride, padding; (void) !fread(&linearisation, sizeof(linearisation), 1, ptr); (void) !fread(&pooling, sizeof(pooling), 1, ptr); + (void) !fread(&stride, sizeof(stride), 1, ptr); + (void) !fread(&padding, sizeof(padding), 1, ptr); kernel->cnn = NULL; kernel->nn = NULL; kernel->activation = IDENTITY; kernel->pooling = pooling; kernel->linearisation = linearisation; + kernel->stride = stride; + kernel->padding = padding; } return kernel; } \ No newline at end of file diff --git a/src/cnn/utils.c b/src/cnn/utils.c index 0ee99ce..9bb864e 100644 --- a/src/cnn/utils.c +++ b/src/cnn/utils.c @@ -45,6 +45,8 @@ bool equals_networks(Network* network1, Network* network2) { for (int i=0; i < network1->size-1; i++) { checkEquals(kernel[i]->activation, "kernel[i]->activation", i); + checkEquals(kernel[i]->stride, "kernel[i]->stride", i); + checkEquals(kernel[i]->padding, "kernel[i]->padding", i); if ((!network1->kernel[i]->cnn ^ !network2->kernel[i]->cnn) || (!network1->kernel[i]->nn ^ !network2->kernel[i]->nn)) { printf(BOLDRED "[ ERROR ]" RESET "network1->kernel[%d] et network1->kernel[%d] diffèrent de type\n", i, i); return false; @@ -129,6 +131,8 @@ Network* copy_network(Network* network) { copyVar(kernel[i]->pooling); copyVar(kernel[i]->activation); copyVar(kernel[i]->linearisation); // 1 + copyVar(kernel[i]->stride); // -1 + copyVar(kernel[i]->padding); // -1 network_cp->kernel[i]->cnn = NULL; network_cp->kernel[i]->nn = NULL; } @@ -136,6 +140,8 @@ Network* copy_network(Network* network) { copyVar(kernel[i]->pooling); copyVar(kernel[i]->activation); copyVar(kernel[i]->linearisation); // 0 + copyVar(kernel[i]->stride); // -1 + copyVar(kernel[i]->padding); // -1 size_input = network->kernel[i]->nn->size_input; size_output = network->kernel[i]->nn->size_output; @@ -188,6 +194,8 @@ Network* copy_network(Network* network) { copyVar(kernel[i]->pooling); copyVar(kernel[i]->activation); copyVar(kernel[i]->linearisation); // 0 + copyVar(kernel[i]->stride); + copyVar(kernel[i]->padding); rows = network->kernel[i]->cnn->rows; k_size = network->kernel[i]->cnn->k_size; diff --git a/test/cnn_structure.c b/test/cnn_structure.c index d1770b6..32bbd56 100644 --- a/test/cnn_structure.c +++ b/test/cnn_structure.c @@ -40,6 +40,8 @@ int main() { 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);