Merge pull request #12291 from cv3d:fix/cuda_pow

modules/cudaarithm/src/cuda/math.cu

@@ -278,20 +278,12 @@ namespace
 {
     template<typename T, bool Signed = numeric_limits<T>::is_signed> struct PowOp : unary_function<T, T>
     {
-        float power;
+        typedef typename LargerType<T, float>::type LargerType;
+        LargerType power;
 
         __device__ __forceinline__ T operator()(T e) const
         {
-            return cudev::saturate_cast<T>(__powf((float)e, power));
-        }
-    };
-    template<typename T> struct PowOp<T, true> : unary_function<T, T>
-    {
-        float power;
-
-        __device__ __forceinline__ T operator()(T e) const
-        {
-            T res = cudev::saturate_cast<T>(__powf((float)e, power));
+            T res = cudev::saturate_cast<T>(__powf(e < 0 ? -e : e, power));
 
             if ((e < 0) && (1 & static_cast<int>(power)))
                 res *= -1;
@@ -299,22 +291,15 @@ namespace
             return res;
         }
     };
-    template<> struct PowOp<float> : unary_function<float, float>
-    {
-        float power;
 
-        __device__ __forceinline__ float operator()(float e) const
-        {
-            return __powf(::fabs(e), power);
-        }
-    };
-    template<> struct PowOp<double> : unary_function<double, double>
+    template<typename T> struct PowOp<T, false> : unary_function<T, T>
     {
-        double power;
+        typedef typename LargerType<T, float>::type LargerType;
+        LargerType power;
 
-        __device__ __forceinline__ double operator()(double e) const
+        __device__ __forceinline__ T operator()(T e) const
         {
-            return ::pow(::fabs(e), power);
+            return cudev::saturate_cast<T>(__powf(e, power));
         }
     };