From 168af40fe9a3cc81c6ee16b3e81f154780c36bdb Mon Sep 17 00:00:00 2001
From: Scheaven <xuepengqiang>
Date: 星期四, 03 六月 2021 15:03:27 +0800
Subject: [PATCH] up new v4
---
lib/detecter_tools/darknet/activation_kernels.cu | 1393 +++++++++++++++++++++++++++++++---------------------------
1 files changed, 745 insertions(+), 648 deletions(-)
diff --git a/lib/detecter_tools/darknet/activation_kernels.cu b/lib/detecter_tools/darknet/activation_kernels.cu
index 62abebd..7a4d06f 100644
--- a/lib/detecter_tools/darknet/activation_kernels.cu
+++ b/lib/detecter_tools/darknet/activation_kernels.cu
@@ -1,649 +1,746 @@
-#include "darknet.h"
-#include <cuda_runtime.h>
-#include <curand.h>
-#include <cublas_v2.h>
-#include <float.h>
-
-#include "activations.h"
-#include "dark_cuda.h"
-
-__device__ float lhtan_activate_kernel(float x)
-{
- if(x < 0) return .001*x;
- if(x > 1) return .001*(x-1) + 1;
- return x;
-}
-__device__ float lhtan_gradient_kernel(float x)
-{
- if(x > 0 && x < 1) return 1;
- return .001;
-}
-
-__device__ float hardtan_activate_kernel(float x)
-{
- if (x < -1) return -1;
- if (x > 1) return 1;
- return x;
-}
-__device__ float linear_activate_kernel(float x){return x;}
-__device__ float logistic_activate_kernel(float x){return 1.f/(1.f + expf(-x));}
-__device__ float loggy_activate_kernel(float x){return 2.f/(1.f + expf(-x)) - 1;}
-__device__ float relu_activate_kernel(float x){return x*(x>0);}
-__device__ float relu6_activate_kernel(float x) { return min_val_cmp(max_val_cmp(x, 0), 6); }
-__device__ float elu_activate_kernel(float x){return (x >= 0)*x + (x < 0)*(expf(x)-1);}
-__device__ float selu_activate_kernel(float x) { return (x >= 0)*1.0507f*x + (x < 0)*1.0507f*1.6732f*(expf(x) - 1); }
-__device__ float relie_activate_kernel(float x){return (x>0) ? x : .01f*x;}
-__device__ float ramp_activate_kernel(float x){return x*(x>0)+.1f*x;}
-__device__ float leaky_activate_kernel(float x){return (x>0) ? x : .1f*x;}
-__device__ float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
-__device__ float gelu_activate_kernel(float x){return (0.5*x*(1 + tanhf(0.797885*x + 0.035677*powf(x, 3))));}
-__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 plse_activate_kernel(float x)
-{
- if(x < -4) return .01f * (x + 4);
- if(x > 4) return .01f * (x - 4) + 1;
- return .125f*x + .5f;
-}
-__device__ float stair_activate_kernel(float x)
-{
- int n = floorf(x);
- if (n%2 == 0) return floorf(x/2.f);
- else return (x - n) + floorf(x/2.f);
-}
-
-
-__device__ float hardtan_gradient_kernel(float x)
-{
- if (x > -1 && x < 1) return 1;
- return 0;
-}
-__device__ float linear_gradient_kernel(float x){return 1;}
-__device__ float logistic_gradient_kernel(float x){return (1-x)*x;}
-__device__ float loggy_gradient_kernel(float x)
-{
- float y = (x+1.F)/2.F;
- return 2*(1-y)*y;
-}
-__device__ float relu_gradient_kernel(float x){return (x>0);}
-__device__ float relu6_gradient_kernel(float x) { return (x > 0 && x < 6); }
-__device__ float elu_gradient_kernel(float x){return (x >= 0) + (x < 0)*(x + 1);}
-__device__ float selu_gradient_kernel(float x) { return (x >= 0)*1.0507f + (x < 0)*(x + 1.0507f*1.6732f); }
-__device__ float relie_gradient_kernel(float x){return (x>0) ? 1 : .01f;}
-__device__ float ramp_gradient_kernel(float x){return (x>0)+.1f;}
-__device__ float leaky_gradient_kernel(float x){return (x>0) ? 1 : .1f;}
-__device__ float tanh_gradient_kernel(float x){return 1-x*x;}
-__device__ float sech_gpu(float x) { return 2 / (expf(x) + expf(-x)); }
-__device__ float gelu_gradient_kernel(float x) {
- const float x3 = powf(x, 3);
- return 0.5*tanhf(0.0356774*x3 + 0.797885*x) + (0.0535161*x3 + 0.398942*x) * powf(sech_gpu(0.0356774*x3 + 0.797885*x), 2) + 0.5;
-}
-__device__ float plse_gradient_kernel(float x){return (x < 0 || x > 1) ? .01f : .125f;}
-__device__ float stair_gradient_kernel(float x)
-{
- if (floor(x) == x) return 0;
- return 1;
-}
-
-__device__ float activate_kernel(float x, ACTIVATION a)
-{
- switch(a){
- case LINEAR:
- return linear_activate_kernel(x);
- case LOGISTIC:
- return logistic_activate_kernel(x);
- case LOGGY:
- return loggy_activate_kernel(x);
- case RELU:
- return relu_activate_kernel(x);
- case RELU6:
- return relu6_activate_kernel(x);
- case ELU:
- return elu_activate_kernel(x);
- case SELU:
- return selu_activate_kernel(x);
- case GELU:
- return gelu_activate_kernel(x);
- case RELIE:
- return relie_activate_kernel(x);
- case RAMP:
- return ramp_activate_kernel(x);
- case LEAKY:
- return leaky_activate_kernel(x);
- case TANH:
- return tanh_activate_kernel(x);
- case PLSE:
- return plse_activate_kernel(x);
- case STAIR:
- return stair_activate_kernel(x);
- case HARDTAN:
- return hardtan_activate_kernel(x);
- case LHTAN:
- return lhtan_activate_kernel(x);
- }
- return 0;
-}
-
-__device__ float gradient_kernel(float x, ACTIVATION a)
-{
- switch (a) {
- case LINEAR:
- return linear_gradient_kernel(x);
- case LOGISTIC:
- return logistic_gradient_kernel(x);
- case LOGGY:
- return loggy_gradient_kernel(x);
- case RELU:
- return relu_gradient_kernel(x);
- case RELU6:
- return relu6_gradient_kernel(x);
- case NORM_CHAN:
- return relu_gradient_kernel(x);
- case ELU:
- return elu_gradient_kernel(x);
- case SELU:
- return selu_gradient_kernel(x);
- case GELU:
- return gelu_gradient_kernel(x);
- case RELIE:
- return relie_gradient_kernel(x);
- case RAMP:
- return ramp_gradient_kernel(x);
- case LEAKY:
- return leaky_gradient_kernel(x);
- case TANH:
- return tanh_gradient_kernel(x);
- case PLSE:
- return plse_gradient_kernel(x);
- case STAIR:
- return stair_gradient_kernel(x);
- case HARDTAN:
- return hardtan_gradient_kernel(x);
- case LHTAN:
- return lhtan_gradient_kernel(x);
- }
- return 0;
-}
-
-__global__ void binary_gradient_array_kernel(float *x, float *dy, int n, int s, BINARY_ACTIVATION a, float *dx)
-{
- int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- int i = id % s;
- int b = id / s;
- float x1 = x[b*s + i];
- float x2 = x[b*s + s / 2 + i];
- if (id < n) {
- float de = dy[id];
- dx[b*s + i] = x2*de;
- dx[b*s + s / 2 + i] = x1*de;
- }
-}
-
-extern "C" void binary_gradient_array_gpu(float *x, float *dx, int n, int size, BINARY_ACTIVATION a, float *y)
-{
- binary_gradient_array_kernel << <cuda_gridsize(n / 2), BLOCK, 0, get_cuda_stream() >> >(x, dx, n / 2, size, a, y);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-__global__ void binary_activate_array_kernel(float *x, int n, int s, BINARY_ACTIVATION a, float *y)
-{
- int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- int i = id % s;
- int b = id / s;
- float x1 = x[b*s + i];
- float x2 = x[b*s + s / 2 + i];
- if (id < n) y[id] = x1*x2;
-}
-
-extern "C" void binary_activate_array_gpu(float *x, int n, int size, BINARY_ACTIVATION a, float *y)
-{
- binary_activate_array_kernel << <cuda_gridsize(n / 2), BLOCK, 0, get_cuda_stream() >> >(x, n / 2, size, a, y);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-__global__ void activate_array_kernel(float *x, int n, ACTIVATION a)
-{
- int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if(i < n) x[i] = activate_kernel(x[i], a);
-}
-
-
-
-__global__ void activate_array_swish_kernel(float *x, int n, float *output_sigmoid_gpu, float *output_gpu)
-{
- int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if (i < n) {
- float x_val = x[i];
- float sigmoid = logistic_activate_kernel(x_val);
- output_sigmoid_gpu[i] = sigmoid;
- output_gpu[i] = x_val * sigmoid;
- }
-}
-
-// https://github.com/digantamisra98/Mish
-__global__ void activate_array_mish_kernel(float *x, int n, float *activation_input, float *output_gpu)
-{
- int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if (i < n) {
- const float MISH_THRESHOLD = 20;
- float x_val = x[i];
- activation_input[i] = x_val; // store value before activation
- //output_gpu[i] = x_val * tanh_activate_kernel(logf(1 + expf(x_val)));
-
- // Pytorch: https://github.com/thomasbrandon/mish-cuda/blob/master/csrc/mish.h#L17-L20
- // TF: https://github.com/tensorflow/addons/blob/093cdfa85d334cbe19a37624c33198f3140109ed/tensorflow_addons/custom_ops/activations/cc/kernels/mish_op.h#L40-L49
- // log1p(x) == log(x + 1)
- output_gpu[i] = x_val * tanh_activate_kernel( softplus_kernel(x_val, MISH_THRESHOLD) );
- }
-}
-
-__global__ void activate_array_leaky_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = leaky_activate_kernel(x[index]);
- }
-}
-
-__global__ void activate_array_selu_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = selu_activate_kernel(x[index]);
- }
-}
-
-__global__ void activate_array_gelu_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = gelu_activate_kernel(x[index]);
- }
-}
-
-__global__ void activate_array_logistic_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = logistic_activate_kernel(x[index]);
- }
-}
-
-__global__ void activate_array_tanh_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = tanh_activate_kernel(x[index]);
- }
-}
-
-__global__ void activate_array_hardtan_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = hardtan_activate_kernel(x[index]);
- }
-}
-
-__global__ void activate_array_relu_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = relu_activate_kernel(x[index]);
- }
-}
-
-__global__ void activate_array_relu6_kernel(float *x, int n)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- x[index] = relu6_activate_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_kernel(float *x, int n, ACTIVATION a, float *delta)
-{
- int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if(i < n) delta[i] *= gradient_kernel(x[i], a);
-}
-
-// https://github.com/BVLC/caffe/blob/04ab089db018a292ae48d51732dd6c66766b36b6/src/caffe/layers/swish_layer.cu#L28-L30
-__global__ void gradient_array_swish_kernel(float *x, int n, float *sigmoid_gpu, float *delta)
-{
- int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if (i < n) {
- float swish = x[i];
- delta[i] *= swish + sigmoid_gpu[i] * (1 - swish); // gradient_kernel(x[i], a);
- }
-}
-
-// https://github.com/digantamisra98/Mish
-__global__ void gradient_array_mish_kernel(int n, float *activation_input_gpu, float *delta)
-{
- int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if (i < n) {
- const float MISH_THRESHOLD = 20.0f;
-
- // implementation from TensorFlow: https://github.com/tensorflow/addons/blob/093cdfa85d334cbe19a37624c33198f3140109ed/tensorflow_addons/custom_ops/activations/cc/kernels/mish_op.h#L66-L80
- // implementation from Pytorch: https://github.com/thomasbrandon/mish-cuda/blob/master/csrc/mish.h#L26-L31
- // log1p(x) == log(x + 1)
- const float inp = activation_input_gpu[i];
- const float sp = softplus_kernel(inp, MISH_THRESHOLD);
- const float grad_sp = 1 - expf(-sp);
- const float tsp = tanh(sp);
- const float grad_tsp = (1 - tsp*tsp) * grad_sp;
- const float grad = inp * grad_tsp + tsp;
- delta[i] *= grad;
-
- //float x = activation_input[i];
- //float d = 2 * expf(x) + expf(2 * x) + 2;
- //float w = 4 * (x + 1) + 4 * expf(2 * x) + expf(3 * x) + expf(x)*(4 * x + 6);
- //float derivative = expf(x) * w / (d * d);
- //delta[i] *= derivative;
- }
-}
-
-__global__ void gradient_array_leaky_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= leaky_gradient_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_selu_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= selu_gradient_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_gelu_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= gelu_gradient_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_logistic_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= logistic_gradient_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_tanh_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= tanh_gradient_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_hardtan_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= hardtan_gradient_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_relu_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= relu_gradient_kernel(x[index]);
- }
-}
-
-__global__ void gradient_array_relu6_kernel(float *x, int n, float *delta)
-{
- int index = blockIdx.x*blockDim.x + threadIdx.x;
- if (index < n) {
- delta[index] *= relu6_gradient_kernel(x[index]);
- }
-}
-
-extern "C" void activate_array_ongpu(float *x, int n, ACTIVATION a)
-{
- const int num_blocks = get_number_of_blocks(n, BLOCK);
- if (a == LINEAR) return;
- else if(a == LEAKY) activate_array_leaky_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else if (a == LOGISTIC) activate_array_logistic_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else if (a == TANH) activate_array_tanh_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else if (a == HARDTAN) activate_array_hardtan_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else if (a == RELU) activate_array_relu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else if (a == RELU6) activate_array_relu6_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else if (a == SELU) activate_array_selu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else if (a == GELU) activate_array_gelu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
- else
- activate_array_kernel<<<cuda_gridsize(n), BLOCK, 0, get_cuda_stream()>>>(x, n, a);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-extern "C" void activate_array_swish_ongpu(float *x, int n, float *output_sigmoid_gpu, float *output_gpu)
-{
- const int num_blocks = get_number_of_blocks(n, BLOCK);
- activate_array_swish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> >(x, n, output_sigmoid_gpu, output_gpu);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-extern "C" void activate_array_mish_ongpu(float *x, int n, float *activation_input_gpu, float *output_gpu)
-{
- const int num_blocks = get_number_of_blocks(n, BLOCK);
- activate_array_mish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> >(x, n, activation_input_gpu, output_gpu);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-extern "C" void gradient_array_ongpu(float *x, int n, ACTIVATION a, float *delta)
-{
- const int num_blocks = get_number_of_blocks(n, BLOCK);
- if (a == LINEAR) return;
- else if (a == LEAKY) gradient_array_leaky_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else if (a == LOGISTIC) gradient_array_logistic_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else if (a == TANH) gradient_array_tanh_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else if (a == HARDTAN) gradient_array_hardtan_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else if (a == RELU) gradient_array_relu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else if (a == RELU6) gradient_array_relu6_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- //else if (a == NORM_CHAN) gradient_array_relu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else if (a == NORM_CHAN_SOFTMAX || a == NORM_CHAN) {
- printf(" Error: should be used custom NORM_CHAN_SOFTMAX-function for gradient \n");
- exit(0);
- }
- else if (a == SELU) gradient_array_selu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else if (a == GELU) gradient_array_gelu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
- else
- gradient_array_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> > (x, n, a, delta);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-
-extern "C" void gradient_array_swish_ongpu(float *x, int n, float *sigmoid_gpu, float *delta)
-{
- const int num_blocks = get_number_of_blocks(n, BLOCK);
- gradient_array_swish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> > (x, n, sigmoid_gpu, delta);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-extern "C" void gradient_array_mish_ongpu(int n, float *activation_input_gpu, float *delta)
-{
- const int num_blocks = get_number_of_blocks(n, BLOCK);
- gradient_array_mish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> > (n, activation_input_gpu, delta);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-
-__global__ void activate_array_normalize_channels_kernel(float *x, int size, int batch, int channels, int wh_step, float *output_gpu)
-{
- int i = blockIdx.x * blockDim.x + threadIdx.x;
-
- int wh_i = i % wh_step;
- int b = i / wh_step;
-
- const float eps = 0.0001;
- if (i < size) {
- float sum = eps;
- int k;
- for (k = 0; k < channels; ++k) {
- float val = x[wh_i + k * wh_step + b*wh_step*channels];
- if (val > 0) sum += val;
- }
- for (k = 0; k < channels; ++k) {
- float val = x[wh_i + k * wh_step + b*wh_step*channels];
- if (val > 0) val = val / sum;
- else val = 0;
- output_gpu[wh_i + k * wh_step + b*wh_step*channels] = val;
- }
- }
-}
-
-extern "C" void activate_array_normalize_channels_ongpu(float *x, int n, int batch, int channels, int wh_step, float *output_gpu)
-{
- // n = w*h*c*batch
- // size = w*h*batch
- int size = n / channels;
-
- const int num_blocks = get_number_of_blocks(size, BLOCK);
-
- activate_array_normalize_channels_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (x, size, batch, channels, wh_step, output_gpu);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-
-
-__global__ void activate_array_normalize_channels_softmax_kernel(float *x, int size, int batch, int channels, int wh_step, float *output_gpu, int use_max_val)
-{
- int i = blockIdx.x * blockDim.x + threadIdx.x;
-
- int wh_i = i % wh_step;
- int b = i / wh_step;
-
- const float eps = 0.0001;
- if (i < size) {
- float sum = eps;
- float max_val = -FLT_MAX;
- int k;
- if (use_max_val) {
- for (k = 0; k < channels; ++k) {
- float val = x[wh_i + k * wh_step + b*wh_step*channels];
- if (val > max_val || k == 0) max_val = val;
- }
- }
- else
- max_val = 0;
-
- for (k = 0; k < channels; ++k) {
- float val = x[wh_i + k * wh_step + b*wh_step*channels];
- sum += expf(val - max_val);
- }
- for (k = 0; k < channels; ++k) {
- float val = x[wh_i + k * wh_step + b*wh_step*channels];
- val = expf(val - max_val) / sum;
- if (isnan(val) || isinf(val)) val = 0;
- output_gpu[wh_i + k * wh_step + b*wh_step*channels] = val;
- }
- }
-}
-
-extern "C" void activate_array_normalize_channels_softmax_ongpu(float *x, int n, int batch, int channels, int wh_step, float *output_gpu, int use_max_val)
-{
- // n = w*h*c*batch
- // size = w*h*batch
- int size = n / channels;
-
- const int num_blocks = get_number_of_blocks(size, BLOCK);
-
- activate_array_normalize_channels_softmax_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (x, size, batch, channels, wh_step, output_gpu, use_max_val);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-
-
-__global__ void gradient_array_normalize_channels_softmax_kernel(float *x, int size, int batch, int channels, int wh_step, float *delta_gpu)
-{
- int i = blockIdx.x * blockDim.x + threadIdx.x;
-
- int wh_i = i % wh_step;
- int b = i / wh_step;
-
- if (i < size) {
- int k;
- /*
- float grad = 0;
- for (k = 0; k < channels; ++k) {
- const int index = wh_i + k * wh_step + b*wh_step*channels;
- float out = x[index];
- float delta = delta_gpu[index];
- grad += out*fabs(delta);
- }
- */
- for (k = 0; k < channels; ++k) {
- const int index = wh_i + k * wh_step + b*wh_step*channels;
- float delta = delta_gpu[index];
- float grad = x[index] * (1 - x[index]);
- delta = delta * grad;
- if (isnan(delta) || isinf(delta)) delta = 0;
- delta_gpu[index] = delta;
- }
- }
-}
-
-extern "C" void gradient_array_normalize_channels_softmax_ongpu(float *output_gpu, int n, int batch, int channels, int wh_step, float *delta_gpu)
-{
- // n = w*h*c*batch
- // size = w*h*batch
- int size = n / channels;
-
- const int num_blocks = get_number_of_blocks(size, BLOCK);
-
- gradient_array_normalize_channels_softmax_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (output_gpu, size, batch, channels, wh_step, delta_gpu);
- CHECK_CUDA(cudaPeekAtLastError());
-}
-
-
-__global__ void gradient_array_normalize_channels_kernel(float *x, int size, int batch, int channels, int wh_step, float *delta_gpu)
-{
- int i = blockIdx.x * blockDim.x + threadIdx.x;
-
- int wh_i = i % wh_step;
- int b = i / wh_step;
-
- if (i < size) {
- int k;
- /*
- float grad = 0;
- for (k = 0; k < channels; ++k) {
- const int index = wh_i + k * wh_step + b*wh_step*channels;
- float out = x[index];
- float delta = delta_gpu[index];
- grad += out*fabs(delta);
- }
- */
- for (k = 0; k < channels; ++k) {
- const int index = wh_i + k * wh_step + b*wh_step*channels;
- if (x[index] > 0) {
- float delta = delta_gpu[index];
- float grad = x[index];
- delta = delta * grad;
- delta_gpu[index] = delta;
- }
- }
- }
-}
-
-extern "C" void gradient_array_normalize_channels_ongpu(float *output_gpu, int n, int batch, int channels, int wh_step, float *delta_gpu)
-{
- // n = w*h*c*batch
- // size = w*h*batch
- int size = n / channels;
-
- const int num_blocks = get_number_of_blocks(size, BLOCK);
-
- gradient_array_normalize_channels_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (output_gpu, size, batch, channels, wh_step, delta_gpu);
- CHECK_CUDA(cudaPeekAtLastError());
+#include "darknet.h"
+#include <cuda_runtime.h>
+#include <curand.h>
+#include <cublas_v2.h>
+#include <float.h>
+
+#include "activations.h"
+#include "dark_cuda.h"
+
+__device__ float lhtan_activate_kernel(float x)
+{
+ if(x < 0) return .001*x;
+ if(x > 1) return .001*(x-1) + 1;
+ return x;
+}
+__device__ float lhtan_gradient_kernel(float x)
+{
+ if(x > 0 && x < 1) return 1;
+ return .001;
+}
+
+__device__ float hardtan_activate_kernel(float x)
+{
+ if (x < -1) return -1;
+ if (x > 1) return 1;
+ return x;
+}
+__device__ float linear_activate_kernel(float x){return x;}
+__device__ float logistic_activate_kernel(float x){return 1.f/(1.f + expf(-x));}
+__device__ float loggy_activate_kernel(float x){return 2.f/(1.f + expf(-x)) - 1;}
+__device__ float relu_activate_kernel(float x){return x*(x>0);}
+__device__ float relu6_activate_kernel(float x) { return min_val_cmp(max_val_cmp(x, 0), 6); }
+__device__ float elu_activate_kernel(float x){return (x >= 0)*x + (x < 0)*(expf(x)-1);}
+__device__ float selu_activate_kernel(float x) { return (x >= 0)*1.0507f*x + (x < 0)*1.0507f*1.6732f*(expf(x) - 1); }
+__device__ float relie_activate_kernel(float x){return (x>0) ? x : .01f*x;}
+__device__ float ramp_activate_kernel(float x){return x*(x>0)+.1f*x;}
+__device__ float leaky_activate_kernel(float x){return (x>0) ? x : .1f*x;}
+__device__ float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
+__device__ float gelu_activate_kernel(float x){return (0.5*x*(1 + tanhf(0.797885*x + 0.035677*powf(x, 3))));}
+__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 log1pf(expf(x));
+ //return logf(expf(x) + 1);
+}
+__device__ float plse_activate_kernel(float x)
+{
+ if(x < -4) return .01f * (x + 4);
+ if(x > 4) return .01f * (x - 4) + 1;
+ return .125f*x + .5f;
+}
+__device__ float stair_activate_kernel(float x)
+{
+ int n = floorf(x);
+ if (n%2 == 0) return floorf(x/2.f);
+ else return (x - n) + floorf(x/2.f);
+}
+
+
+__device__ float hardtan_gradient_kernel(float x)
+{
+ if (x > -1 && x < 1) return 1;
+ return 0;
+}
+__device__ float linear_gradient_kernel(float x){return 1;}
+__device__ float logistic_gradient_kernel(float x){return (1-x)*x;}
+__device__ float loggy_gradient_kernel(float x)
+{
+ float y = (x+1.F)/2.F;
+ return 2*(1-y)*y;
+}
+__device__ float relu_gradient_kernel(float x){return (x>0);}
+__device__ float relu6_gradient_kernel(float x) { return (x > 0 && x < 6); }
+__device__ float elu_gradient_kernel(float x){return (x >= 0) + (x < 0)*(x + 1);}
+__device__ float selu_gradient_kernel(float x) { return (x >= 0)*1.0507f + (x < 0)*(x + 1.0507f*1.6732f); }
+__device__ float relie_gradient_kernel(float x){return (x>0) ? 1 : .01f;}
+__device__ float ramp_gradient_kernel(float x){return (x>0)+.1f;}
+__device__ float leaky_gradient_kernel(float x){return (x>0) ? 1 : .1f;}
+__device__ float tanh_gradient_kernel(float x){return 1-x*x;}
+__device__ float sech_gpu(float x) { return 2 / (expf(x) + expf(-x)); }
+__device__ float gelu_gradient_kernel(float x) {
+ const float x3 = powf(x, 3);
+ return 0.5*tanhf(0.0356774*x3 + 0.797885*x) + (0.0535161*x3 + 0.398942*x) * powf(sech_gpu(0.0356774*x3 + 0.797885*x), 2) + 0.5;
+}
+__device__ float plse_gradient_kernel(float x){return (x < 0 || x > 1) ? .01f : .125f;}
+__device__ float stair_gradient_kernel(float x)
+{
+ if (floorf(x) == x) return 0;
+ return 1;
+}
+
+__device__ float activate_kernel(float x, ACTIVATION a)
+{
+ switch(a){
+ case LINEAR:
+ return linear_activate_kernel(x);
+ case LOGISTIC:
+ return logistic_activate_kernel(x);
+ case LOGGY:
+ return loggy_activate_kernel(x);
+ case RELU:
+ return relu_activate_kernel(x);
+ case RELU6:
+ return relu6_activate_kernel(x);
+ case ELU:
+ return elu_activate_kernel(x);
+ case SELU:
+ return selu_activate_kernel(x);
+ case GELU:
+ return gelu_activate_kernel(x);
+ case RELIE:
+ return relie_activate_kernel(x);
+ case RAMP:
+ return ramp_activate_kernel(x);
+ case LEAKY:
+ return leaky_activate_kernel(x);
+ case TANH:
+ return tanh_activate_kernel(x);
+ case PLSE:
+ return plse_activate_kernel(x);
+ case STAIR:
+ return stair_activate_kernel(x);
+ case HARDTAN:
+ return hardtan_activate_kernel(x);
+ case LHTAN:
+ return lhtan_activate_kernel(x);
+ }
+ return 0;
+}
+
+__device__ float gradient_kernel(float x, ACTIVATION a)
+{
+ switch (a) {
+ case LINEAR:
+ return linear_gradient_kernel(x);
+ case LOGISTIC:
+ return logistic_gradient_kernel(x);
+ case LOGGY:
+ return loggy_gradient_kernel(x);
+ case RELU:
+ return relu_gradient_kernel(x);
+ case RELU6:
+ return relu6_gradient_kernel(x);
+ case NORM_CHAN:
+ return relu_gradient_kernel(x);
+ case ELU:
+ return elu_gradient_kernel(x);
+ case SELU:
+ return selu_gradient_kernel(x);
+ case GELU:
+ return gelu_gradient_kernel(x);
+ case RELIE:
+ return relie_gradient_kernel(x);
+ case RAMP:
+ return ramp_gradient_kernel(x);
+ case LEAKY:
+ return leaky_gradient_kernel(x);
+ case TANH:
+ return tanh_gradient_kernel(x);
+ case PLSE:
+ return plse_gradient_kernel(x);
+ case STAIR:
+ return stair_gradient_kernel(x);
+ case HARDTAN:
+ return hardtan_gradient_kernel(x);
+ case LHTAN:
+ return lhtan_gradient_kernel(x);
+ }
+ return 0;
+}
+
+__global__ void binary_gradient_array_kernel(float *x, float *dy, int n, int s, BINARY_ACTIVATION a, float *dx)
+{
+ int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ int i = id % s;
+ int b = id / s;
+ float x1 = x[b*s + i];
+ float x2 = x[b*s + s / 2 + i];
+ if (id < n) {
+ float de = dy[id];
+ dx[b*s + i] = x2*de;
+ dx[b*s + s / 2 + i] = x1*de;
+ }
+}
+
+extern "C" void binary_gradient_array_gpu(float *x, float *dx, int n, int size, BINARY_ACTIVATION a, float *y)
+{
+ binary_gradient_array_kernel << <cuda_gridsize(n / 2), BLOCK, 0, get_cuda_stream() >> >(x, dx, n / 2, size, a, y);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+__global__ void binary_activate_array_kernel(float *x, int n, int s, BINARY_ACTIVATION a, float *y)
+{
+ int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ int i = id % s;
+ int b = id / s;
+ float x1 = x[b*s + i];
+ float x2 = x[b*s + s / 2 + i];
+ if (id < n) y[id] = x1*x2;
+}
+
+extern "C" void binary_activate_array_gpu(float *x, int n, int size, BINARY_ACTIVATION a, float *y)
+{
+ binary_activate_array_kernel << <cuda_gridsize(n / 2), BLOCK, 0, get_cuda_stream() >> >(x, n / 2, size, a, y);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+__global__ void activate_array_kernel(float *x, int n, ACTIVATION a)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if(i < n) x[i] = activate_kernel(x[i], a);
+}
+
+
+
+__global__ void activate_array_swish_kernel(float *x, int n, float *output_sigmoid_gpu, float *output_gpu)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (i < n) {
+ float x_val = x[i];
+ float sigmoid = logistic_activate_kernel(x_val);
+ if (output_sigmoid_gpu) output_sigmoid_gpu[i] = sigmoid;
+ output_gpu[i] = x_val * sigmoid;
+ }
+}
+
+__device__ float mish_njuffa(float x)
+{
+ float r;
+ float e = expf(x);
+ r = 1.0f / fmaf(fmaf(-0.5f, e, -1.0f), e, -1.0f);
+ r = fmaf(r, x, x);
+ return r;
+}
+
+__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);
+}
+
+__device__ float mish_yashas2(float x)
+{
+ float e = __expf(x);
+ float n = e * e + 2 * e;
+ if (x <= -0.6f)
+ return x * __fdividef(n, n + 2);
+
+ return x - 2 * __fdividef(x, n + 2);
+}
+
+// https://github.com/digantamisra98/Mish
+__global__ void activate_array_mish_kernel(float *x, int n, float *activation_input, float *output_gpu)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (i < n) {
+ const float MISH_THRESHOLD = 20;
+ float x_val = x[i];
+ if (activation_input) activation_input[i] = x_val; // store value before activation
+ //output_gpu[i] = x_val * tanh_activate_kernel(logf(1 + expf(x_val)));
+
+ // Pytorch: https://github.com/thomasbrandon/mish-cuda/blob/master/csrc/mish.h#L17-L20
+ // TF: https://github.com/tensorflow/addons/blob/093cdfa85d334cbe19a37624c33198f3140109ed/tensorflow_addons/custom_ops/activations/cc/kernels/mish_op.h#L40-L49
+ // log1p(x) == log(x + 1)
+ //output_gpu[i] = x_val * tanh_activate_kernel( softplus_kernel(x_val, MISH_THRESHOLD) );
+ output_gpu[i] = mish_yashas2(x_val);
+ //output_gpu[i] = mish_njuffa(x_val);
+ }
+}
+
+__device__ float hard_mish_yashas(float x)
+{
+ if (x > 0)
+ return x;
+ if (x > -2)
+ return x * x / 2 + x;
+ return 0;
+}
+
+__global__ void activate_array_hard_mish_kernel(float *x, int n, float *activation_input, float *output_gpu)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (i < n) {
+
+ float x_val = x[i];
+ if (activation_input) activation_input[i] = x_val; // store value before activation
+ output_gpu[i] = hard_mish_yashas(x_val);
+ }
+}
+__global__ void activate_array_leaky_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = leaky_activate_kernel(x[index]);
+ }
+}
+
+__global__ void activate_array_selu_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = selu_activate_kernel(x[index]);
+ }
+}
+
+__global__ void activate_array_gelu_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = gelu_activate_kernel(x[index]);
+ }
+}
+
+__global__ void activate_array_logistic_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = logistic_activate_kernel(x[index]);
+ }
+}
+
+__global__ void activate_array_tanh_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = tanh_activate_kernel(x[index]);
+ }
+}
+
+__global__ void activate_array_hardtan_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = hardtan_activate_kernel(x[index]);
+ }
+}
+
+__global__ void activate_array_relu_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = relu_activate_kernel(x[index]);
+ }
+}
+
+__global__ void activate_array_relu6_kernel(float *x, int n)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ x[index] = relu6_activate_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_kernel(float *x, int n, ACTIVATION a, float *delta)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if(i < n) delta[i] *= gradient_kernel(x[i], a);
+}
+
+// https://github.com/BVLC/caffe/blob/04ab089db018a292ae48d51732dd6c66766b36b6/src/caffe/layers/swish_layer.cu#L28-L30
+__global__ void gradient_array_swish_kernel(float *x, int n, float *sigmoid_gpu, float *delta)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (i < n) {
+ float swish = x[i];
+ delta[i] *= swish + sigmoid_gpu[i] * (1 - swish); // gradient_kernel(x[i], a);
+ }
+}
+
+// https://github.com/digantamisra98/Mish
+__global__ void gradient_array_mish_kernel(int n, float *activation_input_gpu, float *delta)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (i < n) {
+ const float MISH_THRESHOLD = 20.0f;
+
+ // implementation from TensorFlow: https://github.com/tensorflow/addons/blob/093cdfa85d334cbe19a37624c33198f3140109ed/tensorflow_addons/custom_ops/activations/cc/kernels/mish_op.h#L66-L80
+ // implementation from Pytorch: https://github.com/thomasbrandon/mish-cuda/blob/master/csrc/mish.h#L26-L31
+ // log1p(x) == log(x + 1)
+ const float inp = activation_input_gpu[i];
+ const float sp = softplus_kernel(inp, MISH_THRESHOLD);
+ const float grad_sp = -expm1f(-sp);
+ //const float grad_sp = 1 - expf(-sp);
+ const float tsp = tanh(sp);
+ const float grad_tsp = (1 - tsp*tsp) * grad_sp;
+ const float grad = inp * grad_tsp + tsp;
+ delta[i] *= grad;
+
+ //float x = activation_input[i];
+ //float d = 2 * expf(x) + expf(2 * x) + 2;
+ //float w = 4 * (x + 1) + 4 * expf(2 * x) + expf(3 * x) + expf(x)*(4 * x + 6);
+ //float derivative = expf(x) * w / (d * d);
+ //delta[i] *= derivative;
+ }
+}
+
+__device__ float hard_mish_yashas_grad(float x)
+{
+ if (x > 0)
+ return 1;
+ if (x > -2)
+ return x + 1;
+ return 0;
+}
+
+__global__ void gradient_array_hard_mish_kernel(int n, float *activation_input_gpu, float *delta)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (i < n) {
+
+ const float x = activation_input_gpu[i];
+ delta[i] *= hard_mish_yashas_grad(x);
+ }
+}
+
+__global__ void gradient_array_leaky_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= leaky_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_revleaky_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] /= leaky_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_selu_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= selu_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_gelu_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= gelu_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_logistic_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= logistic_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_tanh_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= tanh_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_hardtan_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= hardtan_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_relu_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= relu_gradient_kernel(x[index]);
+ }
+}
+
+__global__ void gradient_array_relu6_kernel(float *x, int n, float *delta)
+{
+ int index = blockIdx.x*blockDim.x + threadIdx.x;
+ if (index < n) {
+ delta[index] *= relu6_gradient_kernel(x[index]);
+ }
+}
+
+extern "C" void activate_array_ongpu(float *x, int n, ACTIVATION a)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ if (a == LINEAR) return;
+ else if (a == LEAKY || a == REVLEAKY) activate_array_leaky_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else if (a == LOGISTIC) activate_array_logistic_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else if (a == TANH) activate_array_tanh_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else if (a == HARDTAN) activate_array_hardtan_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else if (a == RELU) activate_array_relu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else if (a == RELU6) activate_array_relu6_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else if (a == SELU) activate_array_selu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else if (a == GELU) activate_array_gelu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n);
+ else
+ activate_array_kernel<<<cuda_gridsize(n), BLOCK, 0, get_cuda_stream()>>>(x, n, a);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+extern "C" void activate_array_swish_ongpu(float *x, int n, float *output_sigmoid_gpu, float *output_gpu)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ activate_array_swish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> >(x, n, output_sigmoid_gpu, output_gpu);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+extern "C" void activate_array_mish_ongpu(float *x, int n, float *activation_input_gpu, float *output_gpu)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ activate_array_mish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> >(x, n, activation_input_gpu, output_gpu);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+extern "C" void activate_array_hard_mish_ongpu(float *x, int n, float *activation_input_gpu, float *output_gpu)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ activate_array_hard_mish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> >(x, n, activation_input_gpu, output_gpu);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+extern "C" void gradient_array_ongpu(float *x, int n, ACTIVATION a, float *delta)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ if (a == LINEAR) return;
+ else if (a == LEAKY) gradient_array_leaky_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == REVLEAKY) gradient_array_revleaky_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == LOGISTIC) gradient_array_logistic_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == TANH) gradient_array_tanh_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == HARDTAN) gradient_array_hardtan_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == RELU) gradient_array_relu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == RELU6) gradient_array_relu6_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ //else if (a == NORM_CHAN) gradient_array_relu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == NORM_CHAN_SOFTMAX || a == NORM_CHAN) {
+ printf(" Error: should be used custom NORM_CHAN_SOFTMAX-function for gradient \n");
+ exit(0);
+ }
+ else if (a == SELU) gradient_array_selu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else if (a == GELU) gradient_array_gelu_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(x, n, delta);
+ else
+ gradient_array_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> > (x, n, a, delta);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+
+extern "C" void gradient_array_swish_ongpu(float *x, int n, float *sigmoid_gpu, float *delta)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ gradient_array_swish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> > (x, n, sigmoid_gpu, delta);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+extern "C" void gradient_array_mish_ongpu(int n, float *activation_input_gpu, float *delta)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ gradient_array_mish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> > (n, activation_input_gpu, delta);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+extern "C" void gradient_array_hard_mish_ongpu(int n, float *activation_input_gpu, float *delta)
+{
+ const int num_blocks = get_number_of_blocks(n, BLOCK);
+ gradient_array_hard_mish_kernel << <cuda_gridsize(n), BLOCK, 0, get_cuda_stream() >> > (n, activation_input_gpu, delta);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+
+__global__ void activate_array_normalize_channels_kernel(float *x, int size, int batch, int channels, int wh_step, float *output_gpu)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ int wh_i = i % wh_step;
+ int b = i / wh_step;
+
+ const float eps = 0.0001;
+ if (i < size) {
+ float sum = eps;
+ int k;
+ for (k = 0; k < channels; ++k) {
+ float val = x[wh_i + k * wh_step + b*wh_step*channels];
+ if (val > 0) sum += val;
+ }
+ for (k = 0; k < channels; ++k) {
+ float val = x[wh_i + k * wh_step + b*wh_step*channels];
+ if (val > 0) val = val / sum;
+ else val = 0;
+ output_gpu[wh_i + k * wh_step + b*wh_step*channels] = val;
+ }
+ }
+}
+
+extern "C" void activate_array_normalize_channels_ongpu(float *x, int n, int batch, int channels, int wh_step, float *output_gpu)
+{
+ // n = w*h*c*batch
+ // size = w*h*batch
+ int size = n / channels;
+
+ const int num_blocks = get_number_of_blocks(size, BLOCK);
+
+ activate_array_normalize_channels_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (x, size, batch, channels, wh_step, output_gpu);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+
+
+__global__ void activate_array_normalize_channels_softmax_kernel(float *x, int size, int batch, int channels, int wh_step, float *output_gpu, int use_max_val)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ int wh_i = i % wh_step;
+ int b = i / wh_step;
+
+ const float eps = 0.0001;
+ if (i < size) {
+ float sum = eps;
+ float max_val = -FLT_MAX;
+ int k;
+ if (use_max_val) {
+ for (k = 0; k < channels; ++k) {
+ float val = x[wh_i + k * wh_step + b*wh_step*channels];
+ if (val > max_val || k == 0) max_val = val;
+ }
+ }
+ else
+ max_val = 0;
+
+ for (k = 0; k < channels; ++k) {
+ float val = x[wh_i + k * wh_step + b*wh_step*channels];
+ sum += expf(val - max_val);
+ }
+ for (k = 0; k < channels; ++k) {
+ float val = x[wh_i + k * wh_step + b*wh_step*channels];
+ val = expf(val - max_val) / sum;
+ if (isnan(val) || isinf(val)) val = 0;
+ output_gpu[wh_i + k * wh_step + b*wh_step*channels] = val;
+ }
+ }
+}
+
+extern "C" void activate_array_normalize_channels_softmax_ongpu(float *x, int n, int batch, int channels, int wh_step, float *output_gpu, int use_max_val)
+{
+ // n = w*h*c*batch
+ // size = w*h*batch
+ int size = n / channels;
+
+ const int num_blocks = get_number_of_blocks(size, BLOCK);
+
+ activate_array_normalize_channels_softmax_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (x, size, batch, channels, wh_step, output_gpu, use_max_val);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+
+
+__global__ void gradient_array_normalize_channels_softmax_kernel(float *x, int size, int batch, int channels, int wh_step, float *delta_gpu)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ int wh_i = i % wh_step;
+ int b = i / wh_step;
+
+ if (i < size) {
+ int k;
+ /*
+ float grad = 0;
+ for (k = 0; k < channels; ++k) {
+ const int index = wh_i + k * wh_step + b*wh_step*channels;
+ float out = x[index];
+ float delta = delta_gpu[index];
+ grad += out*fabs(delta);
+ }
+ */
+ for (k = 0; k < channels; ++k) {
+ const int index = wh_i + k * wh_step + b*wh_step*channels;
+ float delta = delta_gpu[index];
+ float grad = x[index] * (1 - x[index]);
+ delta = delta * grad;
+ if (isnan(delta) || isinf(delta)) delta = 0;
+ delta_gpu[index] = delta;
+ }
+ }
+}
+
+extern "C" void gradient_array_normalize_channels_softmax_ongpu(float *output_gpu, int n, int batch, int channels, int wh_step, float *delta_gpu)
+{
+ // n = w*h*c*batch
+ // size = w*h*batch
+ int size = n / channels;
+
+ const int num_blocks = get_number_of_blocks(size, BLOCK);
+
+ gradient_array_normalize_channels_softmax_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (output_gpu, size, batch, channels, wh_step, delta_gpu);
+ CHECK_CUDA(cudaPeekAtLastError());
+}
+
+
+__global__ void gradient_array_normalize_channels_kernel(float *x, int size, int batch, int channels, int wh_step, float *delta_gpu)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ int wh_i = i % wh_step;
+ int b = i / wh_step;
+
+ if (i < size) {
+ int k;
+ /*
+ float grad = 0;
+ for (k = 0; k < channels; ++k) {
+ const int index = wh_i + k * wh_step + b*wh_step*channels;
+ float out = x[index];
+ float delta = delta_gpu[index];
+ grad += out*fabs(delta);
+ }
+ */
+ for (k = 0; k < channels; ++k) {
+ const int index = wh_i + k * wh_step + b*wh_step*channels;
+ if (x[index] > 0) {
+ float delta = delta_gpu[index];
+ float grad = x[index];
+ delta = delta * grad;
+ delta_gpu[index] = delta;
+ }
+ }
+ }
+}
+
+extern "C" void gradient_array_normalize_channels_ongpu(float *output_gpu, int n, int batch, int channels, int wh_step, float *delta_gpu)
+{
+ // n = w*h*c*batch
+ // size = w*h*batch
+ int size = n / channels;
+
+ const int num_blocks = get_number_of_blocks(size, BLOCK);
+
+ gradient_array_normalize_channels_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> > (output_gpu, size, batch, channels, wh_step, delta_gpu);
+ CHECK_CUDA(cudaPeekAtLastError());
}
\ No newline at end of file
--
Gitblit v1.8.0