From d936e5f740d2ead3286a907e5f3310de85e819ae Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Sat, 30 May 2020 16:57:53 +0200 Subject: [PATCH] Add mish_yashas Signed-off-by: Micaela Verucchi --- ...{activation.mish.cu => activation_mish.cu} | 25 ++++++++++++++++--- 1 file changed, 22 insertions(+), 3 deletions(-) rename src/kernels/{activation.mish.cu => activation_mish.cu} (58%) diff --git a/src/kernels/activation.mish.cu b/src/kernels/activation_mish.cu similarity index 58% rename from src/kernels/activation.mish.cu rename to src/kernels/activation_mish.cu index fd55c8a..8900061 100644 --- a/src/kernels/activation.mish.cu +++ b/src/kernels/activation_mish.cu @@ -3,20 +3,39 @@ #define MISH_THRESHOLD 20 -__device__ float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);} -__device__ float softplus_kernel(float x, float threshold = 20) { +__device__ +float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);} + +__device__ +float softplus_kernel(float x, float threshold = 20) { if (x > threshold) return x; // too large else if (x < -threshold) return expf(x); // too small return logf(expf(x) + 1); } + + +__device__ +float mish_yashas(float x) { + float e = __expf(x); + if (x <= -18.0f) + return x * e; + + float n = e * e + 2 * e; + if (x <= -5.0f) + return x * __fdividef(n, n + 2); + + return x - 2 * __fdividef(x, n + 2); +} + // https://github.com/digantamisra98/Mish // https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu __global__ void activation_mish(dnnType *input, dnnType *output, int size) { int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; if (i < size) - output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD)); + // output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD)); + output[i] = mish_yashas(input[i]); } /**