added cv::gpu::pow, ticket #1227

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 4a4df8a..393611e 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -532,6 +532,12 @@ namespace cv
         //! computes exponent of each matrix element (b = e**a)\r
         //! supports only CV_32FC1 type\r
         CV_EXPORTS void exp(const GpuMat& a, GpuMat& b, Stream& stream = Stream::Null());\r
+        \r
+        //! computes power of each matrix element:\r
+        //    (dst(i,j) = pow(     src(i,j) , power), if src.type() is integer\r
+        //    (dst(i,j) = pow(fabs(src(i,j)), power), otherwise\r
+        //! supports all, except depth == CV_64F\r
+        CV_EXPORTS void pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream = Stream::Null());\r
 \r
         //! computes natural logarithm of absolute value of each matrix element: b = log(abs(a))\r
         //! supports only CV_32FC1 type\r

diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu
index 28d4eab..f467361 100644 (file)
--- a/modules/gpu/src/cuda/element_operations.cu
+++ b/modules/gpu/src/cuda/element_operations.cu
@@ -42,6 +42,7 @@
 \r
 #include "opencv2/gpu/device/vecmath.hpp"\r
 #include "opencv2/gpu/device/transform.hpp"\r
+#include "opencv2/gpu/device/limits_gpu.hpp"\r
 #include "opencv2/gpu/device/saturate_cast.hpp"\r
 #include "internal_shared.hpp"\r
 \r
@@ -669,4 +670,63 @@ namespace cv { namespace gpu { namespace mathfunc
     }\r
 \r
     template void subtractCaller<short>(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream);\r
+\r
+\r
+    //////////////////////////////////////////////////////////////////////////\r
+    // pow\r
+    \r
+    template<typename T, bool Signed = device::numeric_limits_gpu<T>::is_signed>\r
+    struct PowOp\r
+    {    \r
+        float power;\r
+        PowOp(float power_) : power(power_) {}\r
+    \r
+        template<typename T>\r
+        __device__ __forceinline__ T operator()(const T& e) const\r
+        {      \r
+            return saturate_cast<T>(__powf((float)e, power));\r
+        }      \r
+    };\r
+\r
+    template<typename T>\r
+    struct PowOp<T, true>\r
+    {\r
+        float power;\r
+        PowOp(float power_) : power(power_) {}\r
+\r
+        __device__ __forceinline__ float operator()(const T& e)\r
+        {\r
+              T res = saturate_cast<T>(__powf((float)e, power));            \r
+            \r
+            if ( (e < 0) && (1 & (int)power) )\r
+                    res *= -1;            \r
+            return res;         \r
+        }\r
+    };\r
+\r
+    template<>\r
+    struct PowOp<float>\r
+    {\r
+        float power;\r
+        PowOp(float power_) : power(power_) {}\r
+\r
+        __device__ __forceinline__ float operator()(const float& e)\r
+        {\r
+            return __powf(fabs(e), power);\r
+        }\r
+    };\r
+\r
+    template<typename T>\r
+    void pow_caller(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream)\r
+    {\r
+        transform((DevMem2D_<T>)src, (DevMem2D_<T>)dst, PowOp<T>(power), stream);\r
+    }   \r
+\r
+    template void pow_caller<uchar>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+    template void pow_caller<schar>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+    template void pow_caller<short>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+    template void pow_caller<ushort>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+    template void pow_caller<int>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+    template void pow_caller<uint>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+    template void pow_caller<float>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
 }}}\r

diff --git a/modules/gpu/src/element_operations.cpp b/modules/gpu/src/element_operations.cpp
index c392df7..4feb819 100644 (file)
--- a/modules/gpu/src/element_operations.cpp
+++ b/modules/gpu/src/element_operations.cpp
@@ -68,6 +68,8 @@ void cv::gpu::max(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(
 void cv::gpu::max(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }\r
 double cv::gpu::threshold(const GpuMat&, GpuMat&, double, double, int, Stream&) {throw_nogpu(); return 0.0;}\r
 \r
+void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&)  { throw_nogpu(); }\r
+\r
 #else\r
 \r
 ////////////////////////////////////////////////////////////////////////\r
@@ -768,4 +770,36 @@ double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh, double
     return thresh;\r
 }\r
 \r
+////////////////////////////////////////////////////////////////////////\r
+// pow\r
+\r
+namespace cv\r
+{\r
+    namespace gpu\r
+    {\r
+        namespace mathfunc\r
+        {\r
+            template<typename T>\r
+            void pow_caller(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+        }\r
+    }\r
+}\r
+\r
+void cv::gpu::pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream)\r
+{    \r
+    CV_Assert( src.depth() != CV_64F );\r
+    dst.create(src.size(), src.type());\r
+\r
+    typedef void (*caller_t)(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
+\r
+    static const caller_t callers[] = \r
+    {\r
+        mathfunc::pow_caller<unsigned char>,  mathfunc::pow_caller<signed char>, \r
+        mathfunc::pow_caller<unsigned short>, mathfunc::pow_caller<short>, \r
+        mathfunc::pow_caller<int>, mathfunc::pow_caller<float>\r
+    };\r
+\r
+    callers[src.depth()](src.reshape(1), (float)power, dst.reshape(1), StreamAccessor::getStream(stream));    \r
+}\r
+\r
 #endif\r

diff --git a/modules/gpu/src/opencv2/gpu/device/limits_gpu.hpp b/modules/gpu/src/opencv2/gpu/device/limits_gpu.hpp
index 28231dc..4baa2f9 100644 (file)
--- a/modules/gpu/src/opencv2/gpu/device/limits_gpu.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/limits_gpu.hpp
@@ -87,6 +87,20 @@ namespace cv { namespace gpu { namespace device
         static const bool is_signed = (char)-1 == -1;\r
     };\r
 \r
+     template<> struct numeric_limits_gpu<signed char>\r
+    {\r
+        typedef char type;\r
+        __device__ __forceinline__ static type min() { return CHAR_MIN; };\r
+        __device__ __forceinline__ static type max() { return CHAR_MAX; };\r
+        __device__ __forceinline__ static type epsilon();\r
+        __device__ __forceinline__ static type round_error();\r
+        __device__ __forceinline__ static type denorm_min();\r
+        __device__ __forceinline__ static type infinity();\r
+        __device__ __forceinline__ static type quiet_NaN();\r
+        __device__ __forceinline__ static type signaling_NaN();\r
+        static const bool is_signed = (signed char)-1 == -1;\r
+    };\r
+\r
     template<> struct numeric_limits_gpu<unsigned char>\r
     {\r
         typedef unsigned char type;\r

diff --git a/modules/gpu/src/stereocsbp.cpp b/modules/gpu/src/stereocsbp.cpp
index 25beb8b..cf02993 100644 (file)
--- a/modules/gpu/src/stereocsbp.cpp
+++ b/modules/gpu/src/stereocsbp.cpp
@@ -107,7 +107,7 @@ void cv::gpu::StereoConstantSpaceBP::estimateRecommendedParams(int width, int he
     levels = (int)::log(static_cast<double>(mm)) * 2 / 3;\r
     if (levels == 0) levels++;\r
 \r
-    nr_plane = (int) ((float) ndisp / pow(2.0, levels + 1));\r
+    nr_plane = (int) ((float) ndisp / std::pow(2.0, levels + 1));\r
 }\r
 \r
 cv::gpu::StereoConstantSpaceBP::StereoConstantSpaceBP(int ndisp_, int iters_, int levels_, int nr_plane_,\r

diff --git a/modules/gpu/test/test_arithm.cpp b/modules/gpu/test/test_arithm.cpp
index 7905dec..0c7494f 100644 (file)
--- a/modules/gpu/test/test_arithm.cpp
+++ b/modules/gpu/test/test_arithm.cpp
@@ -752,6 +752,74 @@ TEST_P(Exp, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, Exp, testing::ValuesIn(devices()));\r
 \r
+\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// pow\r
+\r
+struct Pow : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >\r
+{\r
+    cv::gpu::DeviceInfo devInfo;\r
+    int type;\r
+\r
+    double power;\r
+    cv::Size size;\r
+    cv::Mat mat;\r
+\r
+    cv::Mat dst_gold;\r
+\r
+    virtual void SetUp() \r
+    {        \r
+        devInfo = std::tr1::get<0>(GetParam());\r
+        type = std::tr1::get<1>(GetParam());        \r
+\r
+        cv::gpu::setDevice(devInfo.deviceID());\r
+\r
+        cv::RNG& rng = cvtest::TS::ptr()->get_rng();\r
+\r
+        size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200));\r
+        //size = cv::Size(2, 2);\r
+\r
+        mat = cvtest::randomMat(rng, size, type, 0.0, 100.0, false);        \r
+\r
+        if (mat.depth() == CV_32F)\r
+            power = rng.uniform(1.2f, 3.f);\r
+        else\r
+        {\r
+            int ipower = rng.uniform(2, 8);\r
+            power = (float)ipower;\r
+        }\r
+        cv::pow(mat, power, dst_gold);\r
+    }\r
+};\r
+\r
+TEST_P(Pow, Accuracy) \r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type);\r
+    PRINT_PARAM(size);\r
+    PRINT_PARAM(power);\r
+\r
+    cv::Mat dst;\r
+    \r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat gpu_res;\r
+\r
+        cv::gpu::pow(cv::gpu::GpuMat(mat), power, gpu_res);\r
+\r
+        gpu_res.download(dst);\r
+    );\r
+\r
+    /*std::cout  << mat << std::endl << std::endl;\r
+    std::cout  << dst << std::endl << std::endl;\r
+    std::cout  << dst_gold << std::endl;*/\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1);\r
+}\r
+\r
+INSTANTIATE_TEST_CASE_P(Arithm, Pow, testing::Combine(\r
+                        testing::ValuesIn(devices()),\r
+                        testing::Values(CV_32F, CV_32FC3)));\r
+\r
 ////////////////////////////////////////////////////////////////////////////////\r
 // log\r
 \r

diff --git a/modules/gpu/test/test_main.cpp b/modules/gpu/test/test_main.cpp
index 57a26af..f03fc5a 100644 (file)
--- a/modules/gpu/test/test_main.cpp
+++ b/modules/gpu/test/test_main.cpp
@@ -68,9 +68,13 @@ void print_info()
 #endif\r
 \r
     int deviceCount = cv::gpu::getCudaEnabledDeviceCount();\r
+    int driver;\r
+    cudaDriverGetVersion(&driver);\r
 \r
-    printf("CUDA version: %d\n", CUDART_VERSION);    \r
+    printf("CUDA Driver  version: %d\n", driver);        \r
+    printf("CUDA Runtime version: %d\n", CUDART_VERSION);    \r
     printf("CUDA device count: %d\n\n", deviceCount);\r
+    \r
 \r
     for (int i = 0; i < deviceCount; ++i)\r
     {\r