modules/dnn/src/cuda/math.hpp

   1 // This file is part of OpenCV project.
   2 // It is subject to the license terms in the LICENSE file found in the top-level directory
   3 // of this distribution and at http://opencv.org/license.html.
   4
   5 #ifndef OPENCV_DNN_SRC_CUDA_MATH_HPP
   6 #define OPENCV_DNN_SRC_CUDA_MATH_HPP
   7
   8 #include <cuda_runtime.h>
   9 #include <cuda_fp16.h>
  10
  11 namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace device {
  12
  13     template <class T> __device__ T abs(T val) { return (val < T(0) ? -val : val); }
  14     template <> inline __device__ __half2 abs(__half2 val) {
  15         val.x = abs(val.x);
  16         val.y = abs(val.y);
  17         return val;
  18     }
  19     template <> inline __device__ float abs(float val) { return fabsf(val); }
  20     template <> inline __device__ double abs(double val) { return fabs(val); }
  21
  22     template <class T> __device__ T exp(T val);
  23     template <> inline __device__ __half exp(__half val) { return hexp(val); }
  24     template <> inline __device__ __half2 exp(__half2 val) { return h2exp(val); }
  25     template <> inline __device__ float exp(float val) { return expf(val); }
  26     template <> inline __device__ double exp(double val) { return ::exp(val); }
  27
  28     template <class T> __device__ T expm1(T val);
  29     template <> inline __device__ __half expm1(__half val) { return hexp(val) + __half(1); }
  30     template <> inline __device__ __half2 expm1(__half2 val) { return h2exp(val) + __half2(1, 1); }
  31     template <> inline __device__ float expm1(float val) { return expm1f(val); }
  32     template <> inline __device__ double expm1(double val) { return ::expm1(val); }
  33
  34     template <class T> __device__ T max(T x, T y) { return (x > y ? x : y); }
  35     template <> inline __device__ __half2 max(__half2 a, __half2 b) {
  36         a.x = max(a.x, a.x);
  37         a.y = max(a.y, b.y);
  38         return a;
  39     }
  40     template <> inline __device__ float max(float x, float y) { return fmaxf(x, y); }
  41     template <> inline __device__ double max(double x, double y) { return fmax(x, y); }
  42
  43     template <class T> __device__ T min(T x, T y) { return (x > y ? y : x); }
  44     template <> inline __device__ __half2 min(__half2 a, __half2 b) {
  45         a.x = min(a.x, a.x);
  46         a.y = min(a.y, b.y);
  47         return a;
  48     }
  49     template <> inline __device__ float min(float x, float y) { return fminf(x, y); }
  50     template <> inline __device__ double min(double x, double y) { return fmin(x, y); }
  51
  52     template <class T> __device__ T log1p(T val);
  53     template <> inline __device__ __half log1p(__half val) { return hlog(val) + __half(1); }
  54     template <> inline __device__ __half2 log1p(__half2 val) { return h2log(val) + __half2(1, 1); }
  55     template <> inline __device__ float log1p(float val) { return log1pf(val); }
  56
  57     template <class T> __device__ T log1pexp(T val);
  58     template <> inline __device__ __half log1pexp(__half val) {
  59         if (val <= __half(-4.0))
  60             return exp(val);
  61         else if (val <= __half(8.0))
  62             return log1p(exp(val));
  63         else if (val <= __half(8.7))
  64             return val + exp(-val);
  65         else
  66             return val;
  67     }
  68     template <> inline __device__ __half2 log1pexp(__half2 val) {
  69         val.x = log1pexp(val.x);
  70         val.y = log1pexp(val.y);
  71         return val;
  72     }
  73     template <> inline __device__ float log1pexp(float val) {
  74         if (val <= -20)
  75             return expf(val);
  76         else if (val <= 9.0)
  77             return log1pf(expf(val));
  78         else if (val <= 14.6)
  79             return val + exp(-val);
  80         else
  81             return val;
  82     }
  83     template <> inline __device__ double log1pexp(double val) {
  84         if (val <= -37)
  85             return exp(val);
  86         else if (val <= 18)
  87             return log1p(exp(val));
  88         else if (val <= 33.3)
  89             return val + exp(-val);
  90         else
  91             return val;
  92     }
  93
  94     template <class T> __device__ T tanh(T val);
  95     template <> inline __device__ __half tanh(__half val) { return tanhf(val); }
  96     template <> inline __device__ __half2 tanh(__half2 val) { return __half2(tanh(val.x), tanh(val.y)); }
  97     template <> inline __device__ float tanh(float val) { return tanhf(val); }
  98     template <> inline __device__ double tanh(double val) { return ::tanh(val); }
  99
 100     template <class T> __device__ T pow(T val, T exp);
 101     template <> inline __device__ __half pow(__half val, __half exp) { return powf(val, exp); }
 102     template <> inline __device__ __half2 pow(__half2 val, __half2 exp) { return __half2(pow(val.x, exp.x), pow(val.y, exp.y)); }
 103     template <> inline __device__ float pow(float val, float exp) { return powf(val, exp); }
 104     template <> inline __device__ double pow(double val, double exp) { return ::pow(val, exp); }
 105
 106     template <class T> __device__ T sqrt(T val);
 107     template <> inline __device__ __half sqrt(__half val) { return hsqrt(val); }
 108     template <> inline __device__ __half2 sqrt(__half2 val) { return h2sqrt(val); }
 109     template <> inline __device__ float sqrt(float val) { return sqrtf(val); }
 110     template <> inline __device__ double sqrt(double val) { return ::sqrt(val); }
 111
 112     template <class T> __device__ T rsqrt(T val);
 113     template <> inline __device__ __half rsqrt(__half val) { return hrsqrt(val); }
 114     template <> inline __device__ __half2 rsqrt(__half2 val) { return h2rsqrt(val); }
 115     template <> inline __device__ float rsqrt(float val) { return rsqrtf(val); }
 116     template <> inline __device__ double rsqrt(double val) { return ::rsqrt(val); }
 117
 118     template <class T> __device__ T sigmoid(T val) { return T(1) / (T(1) + exp(-val)); }
 119     template <> inline __device__ __half2 sigmoid(__half2 val) { return __half2(1, 1) / (__half2(1, 1) + exp(__hneg2(val))); }
 120
 121     template <class T> __device__ T clamp(T value, T lower, T upper) { return min(max(value, lower), upper); }
 122
 123 }}}}} /* namespace cv::dnn::cuda4dnn::csl::device */
 124
 125 #endif /* OPENCV_DNN_SRC_CUDA_MATH_HPP */