Fix issues with the network not converging

This commit is contained in:
julienChemillier 2023-05-13 15:39:22 +02:00
parent f316882eeb
commit 3dd2e33fa9
5 changed files with 48 additions and 14 deletions

View File

@ -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]);
}
}

View File

@ -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]);
}
}

View File

@ -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;
}

View File

@ -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;

View File

@ -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);