mirror of
https://github.com/augustin64/projet-tipe
synced 2025-01-23 23:26:25 +01:00
Merge of function in 'src/common/include/utils.h'
This commit is contained in:
parent
46333299bd
commit
3b9ad3db4d
@ -10,11 +10,6 @@
|
|||||||
#include "include/config.h"
|
#include "include/config.h"
|
||||||
|
|
||||||
|
|
||||||
int 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);
|
|
||||||
}
|
|
||||||
|
|
||||||
|
|
||||||
/*
|
/*
|
||||||
* Softmax backward MSE
|
* Softmax backward MSE
|
||||||
*/
|
*/
|
||||||
|
@ -10,11 +10,6 @@
|
|||||||
#include "include/config.h"
|
#include "include/config.h"
|
||||||
|
|
||||||
|
|
||||||
int 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);
|
|
||||||
}
|
|
||||||
|
|
||||||
|
|
||||||
/*
|
/*
|
||||||
* Softmax backward MSE
|
* Softmax backward MSE
|
||||||
*/
|
*/
|
||||||
|
@ -10,9 +10,6 @@
|
|||||||
#ifdef __CUDACC__
|
#ifdef __CUDACC__
|
||||||
__host__ __device__
|
__host__ __device__
|
||||||
#endif
|
#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) {
|
void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, int output_width, int stride, int padding) {
|
||||||
// c'est le kernel de input
|
// c'est le kernel de input
|
||||||
@ -34,7 +31,7 @@ void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, i
|
|||||||
for (int c=-padding; c < max_move; c++) { // colonne du filtre
|
for (int c=-padding; c < max_move; c++) { // colonne du filtre
|
||||||
int x = (stride*j+b);
|
int x = (stride*j+b);
|
||||||
int y = (stride*k+c);
|
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];
|
f += kernel->weights[a][i][b][c]*input[a][stride*j+b][stride*k+c];
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
@ -67,7 +64,7 @@ __global__ void make_convolution_kernel(Kernel_cnn* kernel, float*** input, floa
|
|||||||
for (int c=-padding; c < max_move; c++) {
|
for (int c=-padding; c < max_move; c++) {
|
||||||
int idy_2 = idy*stride+b;
|
int idy_2 = idy*stride+b;
|
||||||
int idz_2 = idz*stride+c;
|
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];
|
f += kernel->weights[a][idx][b][c]*input[a][idy_2][idz_2];
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
@ -10,13 +10,6 @@
|
|||||||
#ifdef __CUDACC__
|
#ifdef __CUDACC__
|
||||||
__host__ __device__
|
__host__ __device__
|
||||||
#endif
|
#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) {
|
void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, int output_width, int stride, int padding) {
|
||||||
// c'est le kernel de input
|
// c'est le kernel de input
|
||||||
@ -38,7 +31,7 @@ void make_convolution_cpu(Kernel_cnn* kernel, float*** input, float*** output, i
|
|||||||
for (int c=-padding; c < max_move; c++) { // colonne du filtre
|
for (int c=-padding; c < max_move; c++) { // colonne du filtre
|
||||||
int x = (stride*j+b);
|
int x = (stride*j+b);
|
||||||
int y = (stride*k+c);
|
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];
|
f += kernel->weights[a][i][b][c]*input[a][stride*j+b][stride*k+c];
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
@ -71,7 +64,7 @@ __global__ void make_convolution_kernel(Kernel_cnn* kernel, float*** input, floa
|
|||||||
for (int c=-padding; c < max_move; c++) {
|
for (int c=-padding; c < max_move; c++) {
|
||||||
int idy_2 = idy*stride+b;
|
int idy_2 = idy*stride+b;
|
||||||
int idz_2 = idz*stride+c;
|
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];
|
f += kernel->weights[a][idx][b][c]*input[a][idy_2][idz_2];
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
@ -3,12 +3,6 @@
|
|||||||
#ifdef __CUDACC__
|
#ifdef __CUDACC__
|
||||||
__host__ __device__
|
__host__ __device__
|
||||||
#endif
|
#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
|
* Effectue la convolution naïvement sur le processeur
|
||||||
|
@ -6,12 +6,6 @@
|
|||||||
#ifdef __CUDACC__
|
#ifdef __CUDACC__
|
||||||
__host__ __device__
|
__host__ __device__
|
||||||
#endif
|
#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
|
* Effectue la propagation d'une convolution avec stride et padding choisis sur le processeur
|
||||||
|
@ -13,9 +13,6 @@
|
|||||||
#ifdef __CUDACC__
|
#ifdef __CUDACC__
|
||||||
__host__ __device__
|
__host__ __device__
|
||||||
#endif
|
#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
|
* Average Pooling
|
||||||
@ -40,7 +37,7 @@ __global__ void make_average_pooling_kernel(float*** input, float*** output, int
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int idy_2 = stride*idy +a;
|
int idy_2 = stride*idy +a;
|
||||||
int idz_2 = stride*idz +b;
|
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];
|
sum += input[idx][idy_2][idz_2];
|
||||||
nb_elements++;
|
nb_elements++;
|
||||||
}
|
}
|
||||||
@ -75,7 +72,7 @@ void make_average_pooling_cpu(float*** input, float*** output, int size, int out
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int j_2 = stride*j +a;
|
int j_2 = stride*j +a;
|
||||||
int k_2 = stride*k +b;
|
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];
|
sum += input[i][j_2][k_2];
|
||||||
nb_elements++;
|
nb_elements++;
|
||||||
}
|
}
|
||||||
@ -125,7 +122,7 @@ __global__ void make_max_pooling_kernel(float*** input, float*** output, int siz
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int idy_2 = stride*idy +a;
|
int idy_2 = stride*idy +a;
|
||||||
int idz_2 = stride*idz +b;
|
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];
|
temp = input[idx][idy_2][idz_2];
|
||||||
m = m > temp ? m : temp; // max(m, temp)
|
m = m > temp ? m : temp; // max(m, temp)
|
||||||
}
|
}
|
||||||
@ -159,7 +156,7 @@ void make_max_pooling_cpu(float*** input, float*** output, int size, int output_
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int j_2 = stride*j +a;
|
int j_2 = stride*j +a;
|
||||||
int k_2 = stride*k +b;
|
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]);
|
m = fmaxf(m, input[i][j_2][k_2]);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
@ -13,9 +13,6 @@
|
|||||||
#ifdef __CUDACC__
|
#ifdef __CUDACC__
|
||||||
__host__ __device__
|
__host__ __device__
|
||||||
#endif
|
#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
|
* Average Pooling
|
||||||
@ -40,7 +37,7 @@ __global__ void make_average_pooling_kernel(float*** input, float*** output, int
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int idy_2 = stride*idy +a;
|
int idy_2 = stride*idy +a;
|
||||||
int idz_2 = stride*idz +b;
|
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];
|
sum += input[idx][idy_2][idz_2];
|
||||||
nb_elements++;
|
nb_elements++;
|
||||||
}
|
}
|
||||||
@ -75,7 +72,7 @@ void make_average_pooling_cpu(float*** input, float*** output, int size, int out
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int j_2 = stride*j +a;
|
int j_2 = stride*j +a;
|
||||||
int k_2 = stride*k +b;
|
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];
|
sum += input[i][j_2][k_2];
|
||||||
nb_elements++;
|
nb_elements++;
|
||||||
}
|
}
|
||||||
@ -125,7 +122,7 @@ __global__ void make_max_pooling_kernel(float*** input, float*** output, int siz
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int idy_2 = stride*idy +a;
|
int idy_2 = stride*idy +a;
|
||||||
int idz_2 = stride*idz +b;
|
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];
|
temp = input[idx][idy_2][idz_2];
|
||||||
m = m > temp ? m : temp; // max(m, temp)
|
m = m > temp ? m : temp; // max(m, temp)
|
||||||
}
|
}
|
||||||
@ -159,7 +156,7 @@ void make_max_pooling_cpu(float*** input, float*** output, int size, int output_
|
|||||||
for (int b=-padding; b < max_move; b++) {
|
for (int b=-padding; b < max_move; b++) {
|
||||||
int j_2 = stride*j +a;
|
int j_2 = stride*j +a;
|
||||||
int k_2 = stride*k +b;
|
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]);
|
m = fmaxf(m, input[i][j_2][k_2]);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
@ -39,6 +39,13 @@ int max(int a, int b);
|
|||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
|
||||||
|
/*
|
||||||
|
* On renvoie true si et seulement si _ et _:
|
||||||
|
* lower_bound <= x < upper_bound
|
||||||
|
* lower_bound <= y < upper_bound
|
||||||
|
*/
|
||||||
|
int not_outside(int x, int y, int lower_bound, int upper_bound);
|
||||||
|
|
||||||
/*
|
/*
|
||||||
* Partie entière supérieure de a/b
|
* Partie entière supérieure de a/b
|
||||||
*/
|
*/
|
||||||
|
@ -26,6 +26,9 @@ int max(int a, int b) {
|
|||||||
}
|
}
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
int 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
|
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);
|
return ((a % b) != 0) ? (a / b + 1) : (a / b);
|
||||||
|
Loading…
Reference in New Issue
Block a user