Luo Tao
8 years ago
commit
62ee11415d
File diff suppressed because it is too large
Load Diff
@ -1,5 +1,3 @@
|
|||||||
if(WITH_AVX)
|
if(WITH_AVX)
|
||||||
cc_library(activation_functions SRCS hl_cpu_functions.cc hl_avx_functions.cc)
|
cc_library(activation_functions SRCS avx_functions.cc)
|
||||||
else()
|
|
||||||
cc_library(activation_functions SRCS hl_cpu_functions.cc)
|
|
||||||
endif()
|
endif()
|
||||||
|
|||||||
@ -0,0 +1,170 @@
|
|||||||
|
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||||
|
|
||||||
|
Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
you may not use this file except in compliance with the License.
|
||||||
|
You may obtain a copy of the License at
|
||||||
|
|
||||||
|
http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
|
||||||
|
Unless required by applicable law or agreed to in writing, software
|
||||||
|
distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
See the License for the specific language governing permissions and
|
||||||
|
limitations under the License. */
|
||||||
|
|
||||||
|
#pragma once
|
||||||
|
#include <math.h>
|
||||||
|
#include "paddle/platform/hostdevice.h"
|
||||||
|
|
||||||
|
#ifdef __AVX__
|
||||||
|
#include <immintrin.h>
|
||||||
|
#endif
|
||||||
|
|
||||||
|
namespace paddle {
|
||||||
|
namespace operators {
|
||||||
|
namespace math {
|
||||||
|
namespace detail {
|
||||||
|
|
||||||
|
#define SIGMOID_THRESHOLD_MIN -40.0
|
||||||
|
#define SIGMOID_THRESHOLD_MAX 13.0
|
||||||
|
#define EXP_MAX_INPUT 40.0
|
||||||
|
|
||||||
|
namespace forward {
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Identity(const T a) {
|
||||||
|
return a;
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Relu(const T a) {
|
||||||
|
return a > static_cast<T>(0.0) ? a : static_cast<T>(0.0);
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Sigmoid(const T a) {
|
||||||
|
const T min = SIGMOID_THRESHOLD_MIN;
|
||||||
|
const T max = SIGMOID_THRESHOLD_MAX;
|
||||||
|
T tmp = (a < min) ? min : ((a > max) ? max : a);
|
||||||
|
return static_cast<T>(1.0) / (static_cast<T>(1.0) + exp(-tmp));
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Tanh(const T a) {
|
||||||
|
T tmp = -2.0 * a;
|
||||||
|
tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp;
|
||||||
|
return (2.0 / (1.0 + exp(tmp))) - 1.0;
|
||||||
|
}
|
||||||
|
|
||||||
|
} // namespace forward
|
||||||
|
|
||||||
|
namespace backward {
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Identity(const T a, const T b) {
|
||||||
|
return a;
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Relu(const T a, const T b) {
|
||||||
|
return a * (b > 0.0 ? 1.0 : 0.0);
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Sigmoid(const T a, const T b) {
|
||||||
|
return a * b * (1.0 - b);
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
DEVICE T Tanh(const T a, const T b) {
|
||||||
|
return a * (1.0 - b * b);
|
||||||
|
}
|
||||||
|
|
||||||
|
} // namespace backward
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
struct Active {
|
||||||
|
typedef T (*Act)(T);
|
||||||
|
typedef T (*ActGrad)(T, T);
|
||||||
|
};
|
||||||
|
|
||||||
|
static DEVICE Active<float>::Act kActFloat[] = {
|
||||||
|
&forward::Sigmoid<float>, &forward::Relu<float>, &forward::Tanh<float>,
|
||||||
|
&forward::Identity<float>};
|
||||||
|
|
||||||
|
static DEVICE Active<float>::ActGrad kActGradFloat[] = {
|
||||||
|
&backward::Sigmoid<float>, &backward::Relu<float>, &backward::Tanh<float>,
|
||||||
|
&backward::Identity<float>};
|
||||||
|
|
||||||
|
static DEVICE Active<double>::Act kActDouble[] = {
|
||||||
|
&forward::Sigmoid<double>, &forward::Relu<double>, &forward::Tanh<double>,
|
||||||
|
&forward::Identity<double>};
|
||||||
|
|
||||||
|
static DEVICE Active<double>::ActGrad kActGradDouble[] = {
|
||||||
|
&backward::Sigmoid<double>, &backward::Relu<double>,
|
||||||
|
&backward::Tanh<double>, &backward::Identity<double>};
|
||||||
|
|
||||||
|
namespace forward {
|
||||||
|
inline DEVICE float activation(float a, int index) {
|
||||||
|
return kActFloat[index](a);
|
||||||
|
}
|
||||||
|
|
||||||
|
inline DEVICE double activation(double a, int index) {
|
||||||
|
return kActDouble[index](a);
|
||||||
|
}
|
||||||
|
|
||||||
|
} // namespace forward
|
||||||
|
|
||||||
|
namespace backward {
|
||||||
|
inline DEVICE float activation(float a, float b, int index) {
|
||||||
|
return kActGradFloat[index](a, b);
|
||||||
|
}
|
||||||
|
|
||||||
|
inline DEVICE double activation(double a, double b, int index) {
|
||||||
|
return kActGradDouble[index](a, b);
|
||||||
|
}
|
||||||
|
} // namespace backward
|
||||||
|
|
||||||
|
#ifdef __AVX__
|
||||||
|
namespace forward {
|
||||||
|
namespace avx {
|
||||||
|
__m256 Relu(const __m256 a);
|
||||||
|
__m256 Sigmoid(const __m256 a);
|
||||||
|
__m256 Tanh(const __m256 a);
|
||||||
|
__m256 Identity(const __m256 a);
|
||||||
|
} // namespace avx
|
||||||
|
} // namespace forward
|
||||||
|
|
||||||
|
namespace backward {
|
||||||
|
namespace avx {
|
||||||
|
__m256 Relu(const __m256 a, const __m256 b);
|
||||||
|
__m256 Sigmoid(const __m256 a, const __m256 b);
|
||||||
|
__m256 Tanh(const __m256 a, const __m256 b);
|
||||||
|
__m256 Identity(const __m256 a, const __m256 b);
|
||||||
|
} // namespace avx
|
||||||
|
} // namespace backward
|
||||||
|
|
||||||
|
static Active<__m256>::Act kActAvx[] = {
|
||||||
|
&forward::avx::Sigmoid, &forward::avx::Relu, &forward::avx::Tanh,
|
||||||
|
&forward::avx::Identity};
|
||||||
|
|
||||||
|
static Active<__m256>::ActGrad kActGradAvx[] = {
|
||||||
|
&backward::avx::Sigmoid, &backward::avx::Relu, &backward::avx::Tanh,
|
||||||
|
&backward::avx::Identity};
|
||||||
|
|
||||||
|
namespace forward {
|
||||||
|
inline __m256 activation(__m256 a, int index) { return kActAvx[index](a); }
|
||||||
|
} // namespace forward
|
||||||
|
|
||||||
|
namespace backward {
|
||||||
|
inline __m256 activation(__m256 a, __m256 b, int index) {
|
||||||
|
return kActGradAvx[index](a, b);
|
||||||
|
}
|
||||||
|
} // namespace backward
|
||||||
|
|
||||||
|
#endif
|
||||||
|
|
||||||
|
} // namespace detail
|
||||||
|
} // namespace math
|
||||||
|
} // namespace operators
|
||||||
|
} // namespace paddle
|
||||||
@ -1,188 +0,0 @@
|
|||||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
|
||||||
|
|
||||||
Licensed under the Apache License, Version 2.0 (the "License");
|
|
||||||
you may not use this file except in compliance with the License.
|
|
||||||
You may obtain a copy of the License at
|
|
||||||
|
|
||||||
http://www.apache.org/licenses/LICENSE-2.0
|
|
||||||
|
|
||||||
Unless required by applicable law or agreed to in writing, software
|
|
||||||
distributed under the License is distributed on an "AS IS" BASIS,
|
|
||||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
||||||
See the License for the specific language governing permissions and
|
|
||||||
limitations under the License. */
|
|
||||||
|
|
||||||
#ifndef HL_ACTIVATION_FUNCTIONS_H_
|
|
||||||
#define HL_ACTIVATION_FUNCTIONS_H_
|
|
||||||
|
|
||||||
#include "hl_functions.h"
|
|
||||||
#include "paddle/operators/math/lstm_compute.h"
|
|
||||||
|
|
||||||
/**
|
|
||||||
* Active functions: sigmoid, relu, tanh and linear.
|
|
||||||
*/
|
|
||||||
#define FLOAT_ACTIVE_FUNCTION \
|
|
||||||
{ \
|
|
||||||
hppl::typef::sigmoid, hppl::typef::relu, hppl::typef::tanh, \
|
|
||||||
hppl::typef::linear \
|
|
||||||
}
|
|
||||||
|
|
||||||
#define DOUBLE_ACTIVE_FUNCTION \
|
|
||||||
{ \
|
|
||||||
hppl::typed::sigmoid, hppl::typed::relu, hppl::typed::tanh, \
|
|
||||||
hppl::typed::linear \
|
|
||||||
}
|
|
||||||
|
|
||||||
#define AVX_ACTIVE_FUNCTION \
|
|
||||||
{ hppl::sigmoid, hppl::relu, hppl::tanh, hppl::linear }
|
|
||||||
|
|
||||||
namespace hppl {
|
|
||||||
|
|
||||||
using activation_mode_t = paddle::operators::math::activation_mode_t;
|
|
||||||
|
|
||||||
/**
|
|
||||||
* Hppl supports sigmoid, relu, tanh, linear active functions
|
|
||||||
* for neural networks' forward and backward activation.
|
|
||||||
*/
|
|
||||||
template <class T>
|
|
||||||
class Active {
|
|
||||||
public:
|
|
||||||
typedef T (*forward)(T);
|
|
||||||
typedef T (*backward)(T, T);
|
|
||||||
};
|
|
||||||
|
|
||||||
template <typename T>
|
|
||||||
struct ForwardActType;
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct ForwardActType<float> {
|
|
||||||
using type = Active<float>::forward;
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct ForwardActType<double> {
|
|
||||||
using type = Active<double>::forward;
|
|
||||||
};
|
|
||||||
|
|
||||||
template <typename T>
|
|
||||||
struct BackwardActType;
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct BackwardActType<float> {
|
|
||||||
using type = Active<float>::backward;
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct BackwardActType<double> {
|
|
||||||
using type = Active<double>::backward;
|
|
||||||
};
|
|
||||||
|
|
||||||
#ifdef __NVCC__
|
|
||||||
namespace gpu {
|
|
||||||
static __device__ Active<float>::forward forward[] = FLOAT_ACTIVE_FUNCTION;
|
|
||||||
static __device__ Active<float>::backward backward[] = FLOAT_ACTIVE_FUNCTION;
|
|
||||||
|
|
||||||
static __device__ Active<double>::forward forward_d[] = DOUBLE_ACTIVE_FUNCTION;
|
|
||||||
static __device__ Active<double>::backward backward_d[] =
|
|
||||||
DOUBLE_ACTIVE_FUNCTION;
|
|
||||||
|
|
||||||
template <typename T>
|
|
||||||
struct ForwardAct {
|
|
||||||
__device__ typename ForwardActType<T>::type operator()(
|
|
||||||
activation_mode_t type);
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct ForwardAct<float> {
|
|
||||||
__device__ ForwardActType<float>::type operator()(activation_mode_t type) {
|
|
||||||
return forward[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct ForwardAct<double> {
|
|
||||||
__device__ ForwardActType<double>::type operator()(activation_mode_t type) {
|
|
||||||
return forward_d[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
template <typename T>
|
|
||||||
struct BackwardAct {
|
|
||||||
__device__ typename BackwardActType<T>::type operator()(
|
|
||||||
activation_mode_t type);
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct BackwardAct<float> {
|
|
||||||
__device__ BackwardActType<float>::type operator()(activation_mode_t type) {
|
|
||||||
return backward[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct BackwardAct<double> {
|
|
||||||
__device__ BackwardActType<double>::type operator()(activation_mode_t type) {
|
|
||||||
return backward_d[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
} // namespace gpu
|
|
||||||
#else
|
|
||||||
namespace cpu {
|
|
||||||
static Active<float>::forward forward[] = FLOAT_ACTIVE_FUNCTION;
|
|
||||||
static Active<float>::backward backward[] = FLOAT_ACTIVE_FUNCTION;
|
|
||||||
|
|
||||||
static Active<double>::forward forward_d[] = DOUBLE_ACTIVE_FUNCTION;
|
|
||||||
static Active<double>::backward backward_d[] = DOUBLE_ACTIVE_FUNCTION;
|
|
||||||
|
|
||||||
template <typename T>
|
|
||||||
struct ForwardAct {
|
|
||||||
typename ForwardActType<T>::type operator()(activation_mode_t type);
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct ForwardAct<float> {
|
|
||||||
ForwardActType<float>::type operator()(activation_mode_t type) {
|
|
||||||
return forward[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct ForwardAct<double> {
|
|
||||||
ForwardActType<double>::type operator()(activation_mode_t type) {
|
|
||||||
return forward_d[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
template <typename T>
|
|
||||||
struct BackwardAct {
|
|
||||||
typename BackwardActType<T>::type operator()(activation_mode_t type);
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct BackwardAct<float> {
|
|
||||||
BackwardActType<float>::type operator()(activation_mode_t type) {
|
|
||||||
return backward[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
template <>
|
|
||||||
struct BackwardAct<double> {
|
|
||||||
BackwardActType<double>::type operator()(activation_mode_t type) {
|
|
||||||
return backward_d[type];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
} // namespace cpu
|
|
||||||
|
|
||||||
#ifdef __AVX__
|
|
||||||
namespace avx {
|
|
||||||
static Active<__m256>::forward forward[] = AVX_ACTIVE_FUNCTION;
|
|
||||||
static Active<__m256>::backward backward[] = AVX_ACTIVE_FUNCTION;
|
|
||||||
} // namespace avx
|
|
||||||
#endif
|
|
||||||
#endif
|
|
||||||
|
|
||||||
} // namespace hppl
|
|
||||||
|
|
||||||
#endif // HL_ACTIVATION_FUNCTIONS_H_
|
|
||||||
@ -1,32 +0,0 @@
|
|||||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
|
||||||
|
|
||||||
Licensed under the Apache License, Version 2.0 (the "License");
|
|
||||||
you may not use this file except in compliance with the License.
|
|
||||||
You may obtain a copy of the License at
|
|
||||||
|
|
||||||
http://www.apache.org/licenses/LICENSE-2.0
|
|
||||||
|
|
||||||
Unless required by applicable law or agreed to in writing, software
|
|
||||||
distributed under the License is distributed on an "AS IS" BASIS,
|
|
||||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
||||||
See the License for the specific language governing permissions and
|
|
||||||
limitations under the License. */
|
|
||||||
|
|
||||||
#ifndef HL_AVX_FUNCTIONS_H_
|
|
||||||
#define HL_AVX_FUNCTIONS_H_
|
|
||||||
|
|
||||||
#include <immintrin.h>
|
|
||||||
|
|
||||||
namespace hppl {
|
|
||||||
__m256 relu(const __m256 a);
|
|
||||||
__m256 sigmoid(const __m256 a);
|
|
||||||
__m256 tanh(const __m256 a);
|
|
||||||
__m256 linear(const __m256 a);
|
|
||||||
|
|
||||||
__m256 relu(const __m256 a, const __m256 b);
|
|
||||||
__m256 sigmoid(const __m256 a, const __m256 b);
|
|
||||||
__m256 tanh(const __m256 a, const __m256 b);
|
|
||||||
__m256 linear(const __m256 a, const __m256 b);
|
|
||||||
} // namespace hppl
|
|
||||||
|
|
||||||
#endif // HL_AVX_FUNCTIONS_H_
|
|
||||||
@ -1,89 +0,0 @@
|
|||||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
|
||||||
|
|
||||||
Licensed under the Apache License, Version 2.0 (the "License");
|
|
||||||
you may not use this file except in compliance with the License.
|
|
||||||
You may obtain a copy of the License at
|
|
||||||
|
|
||||||
http://www.apache.org/licenses/LICENSE-2.0
|
|
||||||
|
|
||||||
Unless required by applicable law or agreed to in writing, software
|
|
||||||
distributed under the License is distributed on an "AS IS" BASIS,
|
|
||||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
||||||
See the License for the specific language governing permissions and
|
|
||||||
limitations under the License. */
|
|
||||||
|
|
||||||
#include <math.h>
|
|
||||||
#include "hl_functions.h"
|
|
||||||
|
|
||||||
namespace hppl {
|
|
||||||
namespace typef {
|
|
||||||
|
|
||||||
float relu(const float a) {
|
|
||||||
return a > static_cast<float>(0.0) ? a : static_cast<float>(0.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
float sigmoid(const float a) {
|
|
||||||
const float min = SIGMOID_THRESHOLD_MIN;
|
|
||||||
const float max = SIGMOID_THRESHOLD_MAX;
|
|
||||||
float tmp = (a < min) ? min : ((a > max) ? max : a);
|
|
||||||
return static_cast<float>(1.0) / (static_cast<float>(1.0) + exp(-tmp));
|
|
||||||
}
|
|
||||||
|
|
||||||
float tanh(const float a) {
|
|
||||||
float tmp = -2.0 * a;
|
|
||||||
tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp;
|
|
||||||
return (2.0 / (1.0 + exp(tmp))) - 1.0;
|
|
||||||
}
|
|
||||||
|
|
||||||
float linear(const float a) { return a; }
|
|
||||||
|
|
||||||
float relu(const float a, const float b) { return a * (b > 0.0 ? 1.0 : 0.0); }
|
|
||||||
|
|
||||||
float sigmoid(const float a, const float b) {
|
|
||||||
return a * b * (static_cast<float>(1) - b);
|
|
||||||
}
|
|
||||||
|
|
||||||
float tanh(const float a, const float b) {
|
|
||||||
return a * (static_cast<float>(1) - b * b);
|
|
||||||
}
|
|
||||||
|
|
||||||
float linear(const float a, const float b) { return a; }
|
|
||||||
|
|
||||||
} // namespace typef
|
|
||||||
|
|
||||||
namespace typed {
|
|
||||||
double relu(const double a) {
|
|
||||||
return a > static_cast<double>(0.0) ? a : static_cast<double>(0.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
double sigmoid(const double a) {
|
|
||||||
const double min = SIGMOID_THRESHOLD_MIN;
|
|
||||||
const double max = SIGMOID_THRESHOLD_MAX;
|
|
||||||
double tmp = (a < min) ? min : ((a > max) ? max : a);
|
|
||||||
return static_cast<double>(1.0) / (static_cast<double>(1.0) + exp(-tmp));
|
|
||||||
}
|
|
||||||
|
|
||||||
double tanh(const double a) {
|
|
||||||
double tmp = -2.0 * a;
|
|
||||||
tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp;
|
|
||||||
return (2.0 / (1.0 + exp(tmp))) - 1.0;
|
|
||||||
}
|
|
||||||
|
|
||||||
double linear(const double a) { return a; }
|
|
||||||
|
|
||||||
double relu(const double a, const double b) {
|
|
||||||
return a * (b > 0.0 ? 1.0 : 0.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
double sigmoid(const double a, const double b) {
|
|
||||||
return a * b * (static_cast<double>(1) - b);
|
|
||||||
}
|
|
||||||
|
|
||||||
double tanh(const double a, const double b) {
|
|
||||||
return a * (static_cast<double>(1) - b * b);
|
|
||||||
}
|
|
||||||
|
|
||||||
double linear(const double a, const double b) { return a; }
|
|
||||||
|
|
||||||
} // namespace typed
|
|
||||||
} // namespace hppl
|
|
||||||
@ -1,71 +0,0 @@
|
|||||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
|
||||||
|
|
||||||
Licensed under the Apache License, Version 2.0 (the "License");
|
|
||||||
you may not use this file except in compliance with the License.
|
|
||||||
You may obtain a copy of the License at
|
|
||||||
|
|
||||||
http://www.apache.org/licenses/LICENSE-2.0
|
|
||||||
|
|
||||||
Unless required by applicable law or agreed to in writing, software
|
|
||||||
distributed under the License is distributed on an "AS IS" BASIS,
|
|
||||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
||||||
See the License for the specific language governing permissions and
|
|
||||||
limitations under the License. */
|
|
||||||
|
|
||||||
#ifndef HL_FUNCTIONS_H_
|
|
||||||
#define HL_FUNCTIONS_H_
|
|
||||||
|
|
||||||
/**
|
|
||||||
* sigmoid threshold maximum
|
|
||||||
*/
|
|
||||||
#define SIGMOID_THRESHOLD_MIN -40.0
|
|
||||||
|
|
||||||
/**
|
|
||||||
* sigmoid threshold minimum
|
|
||||||
*/
|
|
||||||
#define SIGMOID_THRESHOLD_MAX 13.0
|
|
||||||
|
|
||||||
/**
|
|
||||||
* The maximum input value for exp, used to avoid overflow problem.
|
|
||||||
* currently only used for tanh function.
|
|
||||||
*/
|
|
||||||
#define EXP_MAX_INPUT 40.0
|
|
||||||
|
|
||||||
#ifndef __NVCC__
|
|
||||||
namespace hppl {
|
|
||||||
namespace typef {
|
|
||||||
float relu(const float a);
|
|
||||||
float sigmoid(const float a);
|
|
||||||
float tanh(const float a);
|
|
||||||
float linear(const float a);
|
|
||||||
|
|
||||||
float relu(const float a, const float b);
|
|
||||||
float sigmoid(const float a, const float b);
|
|
||||||
float tanh(const float a, const float b);
|
|
||||||
float linear(const float a, const float b);
|
|
||||||
|
|
||||||
} // namespace typef
|
|
||||||
|
|
||||||
namespace typed {
|
|
||||||
double relu(const double a);
|
|
||||||
double sigmoid(const double a);
|
|
||||||
double tanh(const double a);
|
|
||||||
double linear(const double a);
|
|
||||||
|
|
||||||
double relu(const double a, const double b);
|
|
||||||
double sigmoid(const double a, const double b);
|
|
||||||
double tanh(const double a, const double b);
|
|
||||||
double linear(const double a, const double b);
|
|
||||||
} // namespace typed
|
|
||||||
|
|
||||||
} // namespace hppl
|
|
||||||
|
|
||||||
#ifdef __AVX__
|
|
||||||
#include "hl_avx_functions.h"
|
|
||||||
#endif
|
|
||||||
|
|
||||||
#else
|
|
||||||
#include "hl_gpu_functions.h"
|
|
||||||
#endif
|
|
||||||
|
|
||||||
#endif // HL_FUNCTIONS_H_
|
|
||||||
@ -1,93 +0,0 @@
|
|||||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
|
||||||
|
|
||||||
Licensed under the Apache License, Version 2.0 (the "License");
|
|
||||||
you may not use this file except in compliance with the License.
|
|
||||||
You may obtain a copy of the License at
|
|
||||||
|
|
||||||
http://www.apache.org/licenses/LICENSE-2.0
|
|
||||||
|
|
||||||
Unless required by applicable law or agreed to in writing, software
|
|
||||||
distributed under the License is distributed on an "AS IS" BASIS,
|
|
||||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
||||||
See the License for the specific language governing permissions and
|
|
||||||
limitations under the License. */
|
|
||||||
|
|
||||||
#ifndef HL_GPU_FUNCTIONS_CUH_
|
|
||||||
#define HL_GPU_FUNCTIONS_CUH_
|
|
||||||
|
|
||||||
#include "hl_base.h"
|
|
||||||
|
|
||||||
namespace hppl {
|
|
||||||
namespace typef {
|
|
||||||
|
|
||||||
__device__ static float relu(const float a) { return a > 0.0f ? a : 0.0f; }
|
|
||||||
|
|
||||||
__device__ static float sigmoid(const float a) {
|
|
||||||
const float min = SIGMOID_THRESHOLD_MIN;
|
|
||||||
const float max = SIGMOID_THRESHOLD_MAX;
|
|
||||||
float tmp = (a < min) ? min : ((a > max) ? max : a);
|
|
||||||
return __fdividef(1.0f, 1.0f + __expf(-tmp));
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static float tanh(const float a) {
|
|
||||||
float tmp = -2.0 * a;
|
|
||||||
tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp;
|
|
||||||
return __fdividef(2.0f, (1.0f + __expf(-2.0f * tmp))) - 1.0f;
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static float linear(const float a) { return a; }
|
|
||||||
|
|
||||||
__device__ static float relu(const float a, const float b) {
|
|
||||||
return a * (b > 0.0f ? 1.0f : 0.0f);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static float sigmoid(const float a, const float b) {
|
|
||||||
return a * b * (1.0f - b);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static float tanh(const float a, const float b) {
|
|
||||||
return a * (1.0f - b * b);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static float linear(const float a, const float b) { return a; }
|
|
||||||
|
|
||||||
} // namespace typef
|
|
||||||
|
|
||||||
namespace typed {
|
|
||||||
|
|
||||||
__device__ static double relu(const double a) { return a > 0.0 ? a : 0.0; }
|
|
||||||
|
|
||||||
__device__ static double sigmoid(const double a) {
|
|
||||||
const double min = SIGMOID_THRESHOLD_MIN;
|
|
||||||
const double max = SIGMOID_THRESHOLD_MAX;
|
|
||||||
double tmp = (a < min) ? min : ((a > max) ? max : a);
|
|
||||||
return 1.0 / (1.0 + exp(-tmp));
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static double tanh(const double a) {
|
|
||||||
double tmp = -2.0 * a;
|
|
||||||
tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp;
|
|
||||||
return (2.0 / (1.0 + exp(-2.0 * a))) - 1.0;
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static double linear(const double a) { return a; }
|
|
||||||
|
|
||||||
__device__ static double relu(const double a, const double b) {
|
|
||||||
return a * (b > 0.0 ? 1.0 : 0.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static double sigmoid(const double a, const double b) {
|
|
||||||
return a * b * (1 - b);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static double tanh(const double a, const double b) {
|
|
||||||
return a * (1.0 - b * b);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ static double linear(const double a, const double b) { return a; }
|
|
||||||
|
|
||||||
} // namespace typef
|
|
||||||
|
|
||||||
} // namespace hppl
|
|
||||||
|
|
||||||
#endif // HL_GPU_FUNCTIONS_CUH_
|
|
||||||
Loading…
Reference in new issue