Remove duplicated functions

This commit is contained in:
augustin64 2023-03-31 15:16:53 +02:00
parent d6d03162b2
commit 710306a286
3 changed files with 74 additions and 147 deletions

View File

@ -5,26 +5,21 @@
#include "../include/colors.h" #include "../include/colors.h"
#include "../include/utils.h" #include "../include/utils.h"
#include "include/function.h"
#include "include/config.h" #include "include/config.h"
#include "include/function.h"
//* Identity //* Identity
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_identity(float x) { __host__ __device__
return x;
}
__device__ float device_identity_derivative(float x) {
(void)x;
return 1;
}
#endif #endif
float identity(float x) { float identity(float x) {
return x; return x;
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float identity_derivative(float x) { float identity_derivative(float x) {
(void)x; (void)x;
return 1; return 1;
@ -33,20 +28,15 @@ float identity_derivative(float x) {
//* Sigmoid //* Sigmoid
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_sigmoid(float x) { __host__ __device__
return 1/(1 + exp(-x));
}
__device__ float device_sigmoid_derivative(float x) {
float tmp = exp(-x);
return tmp/((1+tmp)*(1+tmp));
}
#endif #endif
float sigmoid(float x) { float sigmoid(float x) {
return 1/(1 + exp(-x)); return 1/(1 + exp(-x));
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float sigmoid_derivative(float x) { float sigmoid_derivative(float x) {
float tmp = exp(-x); float tmp = exp(-x);
return tmp/((1+tmp)*(1+tmp)); return tmp/((1+tmp)*(1+tmp));
@ -55,21 +45,15 @@ float sigmoid_derivative(float x) {
//* RELU //* RELU
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_relu(float x) { __host__ __device__
return fmaxf(0, fminf(x, RELU_CLIP_VALUE));
}
__device__ float device_relu_derivative(float x) {
if (x > 0)
return 1;
return 0;
}
#endif #endif
float relu(float x) { float relu(float x) {
return fmaxf(0, fminf(x, RELU_CLIP_VALUE)); return fmaxf(0, fminf(x, RELU_CLIP_VALUE));
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float relu_derivative(float x) { float relu_derivative(float x) {
if (x > 0) if (x > 0)
return 1; return 1;
@ -79,25 +63,17 @@ float relu_derivative(float x) {
//* Leaky RELU //* Leaky RELU
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_leaky_relu(float x) { __host__ __device__
if (x>0)
return fminf(x, RELU_CLIP_VALUE);
return x*LEAKER;
}
__device__ float device_leaky_relu_derivative(float x) {
if (x > 0)
return 1;
return LEAKER;
}
#endif #endif
float leaky_relu(float x) { float leaky_relu(float x) {
if (x>0) if (x>0)
return fminf(x, RELU_CLIP_VALUE); return fminf(x, RELU_CLIP_VALUE);
return x*LEAKER; return x*LEAKER;
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float leaky_relu_derivative(float x) { float leaky_relu_derivative(float x) {
if (x > 0) if (x > 0)
return 1; return 1;
@ -107,20 +83,15 @@ float leaky_relu_derivative(float x) {
//* Tanh //* Tanh
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_tanh_(float x) { __host__ __device__
return tanh(x);
}
__device__ float device_tanh_derivative(float x) {
float a = tanh(x);
return 1 - a*a;
}
#endif #endif
float tanh_(float x) { float tanh_(float x) {
return tanh(x); return tanh(x);
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float tanh_derivative(float x) { float tanh_derivative(float x) {
float a = tanh(x); float a = tanh(x);
return 1 - a*a; return 1 - a*a;
@ -134,17 +105,17 @@ float tanh_derivative(float x) {
* Définition des pointeurs de fonctions pour CUDA * Définition des pointeurs de fonctions pour CUDA
* voir https://stackoverflow.com/a/15646771 * voir https://stackoverflow.com/a/15646771
*/ */
__device__ funcPtr ptr_sigmoid = device_sigmoid; __device__ funcPtr ptr_sigmoid = sigmoid;
__device__ funcPtr ptr_relu = device_relu; __device__ funcPtr ptr_relu = relu;
__device__ funcPtr ptr_leaky_relu = device_leaky_relu; __device__ funcPtr ptr_leaky_relu = leaky_relu;
__device__ funcPtr ptr_tanh = device_tanh_; __device__ funcPtr ptr_tanh = tanh_;
__device__ funcPtr ptr_identity = device_identity; __device__ funcPtr ptr_identity = identity;
__device__ funcPtr ptr_identity_derivative = device_identity_derivative; __device__ funcPtr ptr_identity_derivative = identity_derivative;
__device__ funcPtr ptr_sigmoid_derivative = device_sigmoid_derivative; __device__ funcPtr ptr_sigmoid_derivative = sigmoid_derivative;
__device__ funcPtr ptr_relu_derivative = device_relu_derivative; __device__ funcPtr ptr_relu_derivative = relu_derivative;
__device__ funcPtr ptr_leaky_relu_derivative = device_leaky_relu_derivative; __device__ funcPtr ptr_leaky_relu_derivative = leaky_relu_derivative;
__device__ funcPtr ptr_tanh_derivative = device_tanh_derivative; __device__ funcPtr ptr_tanh_derivative = tanh_derivative;
#endif #endif

View File

@ -5,26 +5,21 @@
#include "../include/colors.h" #include "../include/colors.h"
#include "../include/utils.h" #include "../include/utils.h"
#include "include/function.h"
#include "include/config.h" #include "include/config.h"
#include "include/function.h"
//* Identity //* Identity
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_identity(float x) { __host__ __device__
return x;
}
__device__ float device_identity_derivative(float x) {
(void)x;
return 1;
}
#endif #endif
float identity(float x) { float identity(float x) {
return x; return x;
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float identity_derivative(float x) { float identity_derivative(float x) {
(void)x; (void)x;
return 1; return 1;
@ -33,20 +28,15 @@ float identity_derivative(float x) {
//* Sigmoid //* Sigmoid
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_sigmoid(float x) { __host__ __device__
return 1/(1 + exp(-x));
}
__device__ float device_sigmoid_derivative(float x) {
float tmp = exp(-x);
return tmp/((1+tmp)*(1+tmp));
}
#endif #endif
float sigmoid(float x) { float sigmoid(float x) {
return 1/(1 + exp(-x)); return 1/(1 + exp(-x));
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float sigmoid_derivative(float x) { float sigmoid_derivative(float x) {
float tmp = exp(-x); float tmp = exp(-x);
return tmp/((1+tmp)*(1+tmp)); return tmp/((1+tmp)*(1+tmp));
@ -55,21 +45,15 @@ float sigmoid_derivative(float x) {
//* RELU //* RELU
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_relu(float x) { __host__ __device__
return fmaxf(0, fminf(x, RELU_CLIP_VALUE));
}
__device__ float device_relu_derivative(float x) {
if (x > 0)
return 1;
return 0;
}
#endif #endif
float relu(float x) { float relu(float x) {
return fmaxf(0, fminf(x, RELU_CLIP_VALUE)); return fmaxf(0, fminf(x, RELU_CLIP_VALUE));
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float relu_derivative(float x) { float relu_derivative(float x) {
if (x > 0) if (x > 0)
return 1; return 1;
@ -79,25 +63,17 @@ float relu_derivative(float x) {
//* Leaky RELU //* Leaky RELU
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_leaky_relu(float x) { __host__ __device__
if (x>0)
return fminf(x, RELU_CLIP_VALUE);
return x*LEAKER;
}
__device__ float device_leaky_relu_derivative(float x) {
if (x > 0)
return 1;
return LEAKER;
}
#endif #endif
float leaky_relu(float x) { float leaky_relu(float x) {
if (x>0) if (x>0)
return fminf(x, RELU_CLIP_VALUE); return fminf(x, RELU_CLIP_VALUE);
return x*LEAKER; return x*LEAKER;
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float leaky_relu_derivative(float x) { float leaky_relu_derivative(float x) {
if (x > 0) if (x > 0)
return 1; return 1;
@ -107,20 +83,15 @@ float leaky_relu_derivative(float x) {
//* Tanh //* Tanh
#ifdef __CUDACC__ #ifdef __CUDACC__
__device__ float device_tanh_(float x) { __host__ __device__
return tanh(x);
}
__device__ float device_tanh_derivative(float x) {
float a = tanh(x);
return 1 - a*a;
}
#endif #endif
float tanh_(float x) { float tanh_(float x) {
return tanh(x); return tanh(x);
} }
#ifdef __CUDACC__
__host__ __device__
#endif
float tanh_derivative(float x) { float tanh_derivative(float x) {
float a = tanh(x); float a = tanh(x);
return 1 - a*a; return 1 - a*a;
@ -134,17 +105,17 @@ float tanh_derivative(float x) {
* Définition des pointeurs de fonctions pour CUDA * Définition des pointeurs de fonctions pour CUDA
* voir https://stackoverflow.com/a/15646771 * voir https://stackoverflow.com/a/15646771
*/ */
__device__ funcPtr ptr_sigmoid = device_sigmoid; __device__ funcPtr ptr_sigmoid = sigmoid;
__device__ funcPtr ptr_relu = device_relu; __device__ funcPtr ptr_relu = relu;
__device__ funcPtr ptr_leaky_relu = device_leaky_relu; __device__ funcPtr ptr_leaky_relu = leaky_relu;
__device__ funcPtr ptr_tanh = device_tanh_; __device__ funcPtr ptr_tanh = tanh_;
__device__ funcPtr ptr_identity = device_identity; __device__ funcPtr ptr_identity = identity;
__device__ funcPtr ptr_identity_derivative = device_identity_derivative; __device__ funcPtr ptr_identity_derivative = identity_derivative;
__device__ funcPtr ptr_sigmoid_derivative = device_sigmoid_derivative; __device__ funcPtr ptr_sigmoid_derivative = sigmoid_derivative;
__device__ funcPtr ptr_relu_derivative = device_relu_derivative; __device__ funcPtr ptr_relu_derivative = relu_derivative;
__device__ funcPtr ptr_leaky_relu_derivative = device_leaky_relu_derivative; __device__ funcPtr ptr_leaky_relu_derivative = leaky_relu_derivative;
__device__ funcPtr ptr_tanh_derivative = device_tanh_derivative; __device__ funcPtr ptr_tanh_derivative = tanh_derivative;
#endif #endif

View File

@ -19,82 +19,67 @@
typedef float (*funcPtr)(float); typedef float (*funcPtr)(float);
//* Identité //* Identité
#ifdef __CUDACC__
__device__ float device_identity(float x);
__device__ float device_identity_derivative(float x);
#endif
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float identity(float x); float identity(float x);
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float identity_derivative(float x); float identity_derivative(float x);
//* Sigmoid //* Sigmoid
#ifdef __CUDACC__
__device__ float device_sigmoid(float x);
__device__ float device_sigmoid_derivative(float x);
#endif
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float sigmoid(float x); float sigmoid(float x);
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float sigmoid_derivative(float x); float sigmoid_derivative(float x);
//* RELU //* RELU
#ifdef __CUDACC__
__device__ float device_relu(float x);
__device__ float device_relu_derivative(float x);
#endif
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float relu(float x); float relu(float x);
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float relu_derivative(float x); float relu_derivative(float x);
//* Leaky RELU //* Leaky RELU
#ifdef __CUDACC__
__device__ float device_leaky_relu(float x);
__device__ float device_leaky_relu_derivative(float x);
#endif
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float leaky_relu(float x); float leaky_relu(float x);
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float leaky_relu_derivative(float x); float leaky_relu_derivative(float x);
//* Tanh //* Tanh
#ifdef __CUDACC__
__device__ float device_tanh_(float x);
__device__ float device_tanh_derivative(float x);
#endif
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float tanh_(float x); float tanh_(float x);
#ifdef __CUDACC__ #ifdef __CUDACC__
extern "C" extern "C"
__host__ __device__
#endif #endif
float tanh_derivative(float x); float tanh_derivative(float x);