added is_signed into numeric_limits_gpu, fixed incorrect min max finding for floating...

diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp
index c7ca547..48bed39 100644 (file)
--- a/modules/gpu/src/arithm.cpp
+++ b/modules/gpu/src/arithm.cpp
@@ -524,20 +524,20 @@ void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal, const Gp
 \r
     typedef void (*Caller)(const DevMem2D, double*, double*, PtrStep);\r
     static const Caller callers[2][7] = \r
-    { { min_max_multipass_caller<unsigned char>, min_max_multipass_caller<signed char>, \r
-        min_max_multipass_caller<unsigned short>, min_max_multipass_caller<signed short>, \r
+    { { min_max_multipass_caller<unsigned char>, min_max_multipass_caller<char>, \r
+        min_max_multipass_caller<unsigned short>, min_max_multipass_caller<short>, \r
         min_max_multipass_caller<int>, min_max_multipass_caller<float>, 0 },\r
-      { min_max_caller<unsigned char>, min_max_caller<signed char>, \r
-        min_max_caller<unsigned short>, min_max_caller<signed short>, \r
+      { min_max_caller<unsigned char>, min_max_caller<char>, \r
+        min_max_caller<unsigned short>, min_max_caller<short>, \r
         min_max_caller<int>, min_max_caller<float>, min_max_caller<double> } };\r
 \r
     typedef void (*MaskedCaller)(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     static const MaskedCaller masked_callers[2][7] = \r
-    { { min_max_mask_multipass_caller<unsigned char>, min_max_mask_multipass_caller<signed char>, \r
-        min_max_mask_multipass_caller<unsigned short>, min_max_mask_multipass_caller<signed short>, \r
+    { { min_max_mask_multipass_caller<unsigned char>, min_max_mask_multipass_caller<char>, \r
+        min_max_mask_multipass_caller<unsigned short>, min_max_mask_multipass_caller<short>, \r
         min_max_mask_multipass_caller<int>, min_max_mask_multipass_caller<float>, 0 },\r
-      { min_max_mask_caller<unsigned char>, min_max_mask_caller<signed char>, \r
-        min_max_mask_caller<unsigned short>, min_max_mask_caller<signed short>, \r
+      { min_max_mask_caller<unsigned char>, min_max_mask_caller<char>, \r
+        min_max_mask_caller<unsigned short>, min_max_mask_caller<short>, \r
         min_max_mask_caller<int>, min_max_mask_caller<float>, \r
         min_max_mask_caller<double> } };\r
 \r
@@ -615,9 +615,9 @@ void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point
         switch (src.type())\r
         {\r
         case CV_8U: min_max_loc_caller<unsigned char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
-        case CV_8S: min_max_loc_caller<signed char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
+        case CV_8S: min_max_loc_caller<char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         case CV_16U: min_max_loc_caller<unsigned short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
-        case CV_16S: min_max_loc_caller<signed short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
+        case CV_16S: min_max_loc_caller<short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         case CV_32S: min_max_loc_caller<int>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         case CV_32F: min_max_loc_caller<float>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         case CV_64F: \r
@@ -634,9 +634,9 @@ void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point
         switch (src.type())\r
         {\r
         case CV_8U: min_max_loc_multipass_caller<unsigned char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
-        case CV_8S: min_max_loc_multipass_caller<signed char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
+        case CV_8S: min_max_loc_multipass_caller<char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         case CV_16U: min_max_loc_multipass_caller<unsigned short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
-        case CV_16S: min_max_loc_multipass_caller<signed short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
+        case CV_16S: min_max_loc_multipass_caller<short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         case CV_32S: min_max_loc_multipass_caller<int>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         case CV_32F: min_max_loc_multipass_caller<float>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
         default: CV_Error(CV_StsBadArg, "minMaxLoc: unsupported type");\r
@@ -683,9 +683,9 @@ int cv::gpu::countNonZero(const GpuMat& src, GpuMat& buf)
         switch (src.type())\r
         {\r
         case CV_8U: return count_non_zero_caller<unsigned char>(src, buf);\r
-        case CV_8S: return count_non_zero_caller<signed char>(src, buf);\r
+        case CV_8S: return count_non_zero_caller<char>(src, buf);\r
         case CV_16U: return count_non_zero_caller<unsigned short>(src, buf);\r
-        case CV_16S: return count_non_zero_caller<signed short>(src, buf);\r
+        case CV_16S: return count_non_zero_caller<short>(src, buf);\r
         case CV_32S: return count_non_zero_caller<int>(src, buf);\r
         case CV_32F: return count_non_zero_caller<float>(src, buf);\r
         case CV_64F: \r
@@ -698,9 +698,9 @@ int cv::gpu::countNonZero(const GpuMat& src, GpuMat& buf)
         switch (src.type())\r
         {\r
         case CV_8U: return count_non_zero_multipass_caller<unsigned char>(src, buf);\r
-        case CV_8S: return count_non_zero_multipass_caller<signed char>(src, buf);\r
+        case CV_8S: return count_non_zero_multipass_caller<char>(src, buf);\r
         case CV_16U: return count_non_zero_multipass_caller<unsigned short>(src, buf);\r
-        case CV_16S: return count_non_zero_multipass_caller<signed short>(src, buf);\r
+        case CV_16S: return count_non_zero_multipass_caller<short>(src, buf);\r
         case CV_32S: return count_non_zero_multipass_caller<int>(src, buf);\r
         case CV_32F: return count_non_zero_multipass_caller<float>(src, buf);\r
         }\r

diff --git a/modules/gpu/src/cuda/limits_gpu.hpp b/modules/gpu/src/cuda/limits_gpu.hpp
index 83a0743..fcf4bdc 100644 (file)
--- a/modules/gpu/src/cuda/limits_gpu.hpp
+++ b/modules/gpu/src/cuda/limits_gpu.hpp
@@ -58,6 +58,7 @@ namespace cv
                 __device__ static type infinity() { return type(); }\r
                 __device__ static type quiet_NaN() { return type(); }\r
                 __device__ static type signaling_NaN() { return T(); }\r
+                static const bool is_signed;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<bool>\r
@@ -71,6 +72,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = false;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<char>\r
@@ -84,6 +86,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = (char)-1 == -1;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<unsigned char>\r
@@ -97,6 +100,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = false;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<short>\r
@@ -110,6 +114,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = true;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<unsigned short>\r
@@ -123,6 +128,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = false;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<int>\r
@@ -136,6 +142,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = true;\r
             };\r
 \r
 \r
@@ -150,6 +157,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = false;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<long>\r
@@ -163,6 +171,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = true;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<unsigned long>\r
@@ -176,6 +185,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = false;\r
             };\r
                         \r
             template<> struct numeric_limits_gpu<float>\r
@@ -189,6 +199,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = true;\r
             };\r
 \r
             template<> struct numeric_limits_gpu<double>\r
@@ -202,6 +213,7 @@ namespace cv
                 __device__ static type infinity();\r
                 __device__ static type quiet_NaN();\r
                 __device__ static type signaling_NaN();\r
+                static const bool is_signed = true;\r
             };            \r
         }\r
     }\r

diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu
index d4ecead..b00c94f 100644 (file)
--- a/modules/gpu/src/cuda/mathfunc.cu
+++ b/modules/gpu/src/cuda/mathfunc.cu
@@ -405,9 +405,9 @@ namespace cv { namespace gpu { namespace mathfunc
     // appropriate type (32 bits minimum)\r
     template <typename T> struct MinMaxTypeTraits {};\r
     template <> struct MinMaxTypeTraits<unsigned char> { typedef int best_type; };\r
-    template <> struct MinMaxTypeTraits<signed char> { typedef int best_type; };\r
+    template <> struct MinMaxTypeTraits<char> { typedef int best_type; };\r
     template <> struct MinMaxTypeTraits<unsigned short> { typedef int best_type; };\r
-    template <> struct MinMaxTypeTraits<signed short> { typedef int best_type; };\r
+    template <> struct MinMaxTypeTraits<short> { typedef int best_type; };\r
     template <> struct MinMaxTypeTraits<int> { typedef int best_type; };\r
     template <> struct MinMaxTypeTraits<float> { typedef float best_type; };\r
     template <> struct MinMaxTypeTraits<double> { typedef double best_type; };\r
@@ -492,7 +492,7 @@ namespace cv { namespace gpu { namespace mathfunc
         unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
 \r
         T mymin = numeric_limits_gpu<T>::max();\r
-        T mymax = numeric_limits_gpu<T>::min();\r
+        T mymax = numeric_limits_gpu<T>::is_signed ? -numeric_limits_gpu<T>::max() : numeric_limits_gpu<T>::min();\r
         unsigned int y_end = min(y0 + (ctheight - 1) * blockDim.y + 1, src.rows);\r
         unsigned int x_end = min(x0 + (ctwidth - 1) * blockDim.x + 1, src.cols);\r
         for (unsigned int y = y0; y < y_end; y += blockDim.y)\r
@@ -584,9 +584,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }  \r
 \r
     template void min_max_mask_caller<unsigned char>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
-    template void min_max_mask_caller<signed char>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
+    template void min_max_mask_caller<char>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     template void min_max_mask_caller<unsigned short>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
-    template void min_max_mask_caller<signed short>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
+    template void min_max_mask_caller<short>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     template void min_max_mask_caller<int>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     template void min_max_mask_caller<float>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     template void min_max_mask_caller<double>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
@@ -613,9 +613,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }  \r
 \r
     template void min_max_caller<unsigned char>(const DevMem2D, double*, double*, PtrStep);\r
-    template void min_max_caller<signed char>(const DevMem2D, double*, double*, PtrStep);\r
+    template void min_max_caller<char>(const DevMem2D, double*, double*, PtrStep);\r
     template void min_max_caller<unsigned short>(const DevMem2D, double*, double*, PtrStep);\r
-    template void min_max_caller<signed short>(const DevMem2D, double*, double*, PtrStep);\r
+    template void min_max_caller<short>(const DevMem2D, double*, double*, PtrStep);\r
     template void min_max_caller<int>(const DevMem2D, double*, double*, PtrStep);\r
     template void min_max_caller<float>(const DevMem2D, double*,double*, PtrStep);\r
     template void min_max_caller<double>(const DevMem2D, double*, double*, PtrStep);\r
@@ -668,9 +668,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }\r
 \r
     template void min_max_mask_multipass_caller<unsigned char>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
-    template void min_max_mask_multipass_caller<signed char>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
+    template void min_max_mask_multipass_caller<char>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     template void min_max_mask_multipass_caller<unsigned short>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
-    template void min_max_mask_multipass_caller<signed short>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
+    template void min_max_mask_multipass_caller<short>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     template void min_max_mask_multipass_caller<int>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
     template void min_max_mask_multipass_caller<float>(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
 \r
@@ -697,9 +697,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }\r
 \r
     template void min_max_multipass_caller<unsigned char>(const DevMem2D, double*, double*, PtrStep);\r
-    template void min_max_multipass_caller<signed char>(const DevMem2D, double*, double*, PtrStep);\r
+    template void min_max_multipass_caller<char>(const DevMem2D, double*, double*, PtrStep);\r
     template void min_max_multipass_caller<unsigned short>(const DevMem2D, double*, double*, PtrStep);\r
-    template void min_max_multipass_caller<signed short>(const DevMem2D, double*, double*, PtrStep);\r
+    template void min_max_multipass_caller<short>(const DevMem2D, double*, double*, PtrStep);\r
     template void min_max_multipass_caller<int>(const DevMem2D, double*, double*, PtrStep);\r
     template void min_max_multipass_caller<float>(const DevMem2D, double*, double*, PtrStep);\r
 \r
@@ -802,10 +802,10 @@ namespace cv { namespace gpu { namespace mathfunc
         unsigned int y0 = blockIdx.y * blockDim.y * ctheight + threadIdx.y;\r
         unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
 \r
-        T val = ((const T*)src.ptr(0))[0];\r
-        T mymin = val, mymax = val; \r
-        unsigned int myminloc = 0, mymaxloc = 0;\r
-\r
+        T mymin = numeric_limits_gpu<T>::max();\r
+        T mymax = numeric_limits_gpu<T>::is_signed ? -numeric_limits_gpu<T>::max() : numeric_limits_gpu<T>::min(); \r
+        unsigned int myminloc = 0;\r
+        unsigned int mymaxloc = 0;\r
         unsigned int y_end = min(y0 + (ctheight - 1) * blockDim.y + 1, src.rows);\r
         unsigned int x_end = min(x0 + (ctwidth - 1) * blockDim.x + 1, src.cols);\r
 \r
@@ -814,13 +814,13 @@ namespace cv { namespace gpu { namespace mathfunc
             const T* ptr = (const T*)src.ptr(y);\r
             for (unsigned int x = x0; x < x_end; x += blockDim.x)\r
             {\r
-                val = ptr[x];\r
-                if (val < mymin) \r
+                T val = ptr[x];\r
+                if (val <= mymin) \r
                 { \r
                     mymin = val; \r
                     myminloc = y * src.cols + x; \r
                 }\r
-                else if (val > mymax)\r
+                if (val >= mymax)\r
                 {\r
                     mymax = val; \r
                     mymaxloc = y * src.cols + x; \r
@@ -916,9 +916,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }\r
 \r
     template void min_max_loc_caller<unsigned char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
-    template void min_max_loc_caller<signed char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_caller<char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
     template void min_max_loc_caller<unsigned short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
-    template void min_max_loc_caller<signed short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_caller<short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
     template void min_max_loc_caller<int>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
     template void min_max_loc_caller<float>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
     template void min_max_loc_caller<double>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
@@ -987,9 +987,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }\r
 \r
     template void min_max_loc_multipass_caller<unsigned char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
-    template void min_max_loc_multipass_caller<signed char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_multipass_caller<char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
     template void min_max_loc_multipass_caller<unsigned short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
-    template void min_max_loc_multipass_caller<signed short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_multipass_caller<short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
     template void min_max_loc_multipass_caller<int>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
     template void min_max_loc_multipass_caller<float>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
 \r
@@ -1126,9 +1126,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }  \r
 \r
     template int count_non_zero_caller<unsigned char>(const DevMem2D, PtrStep);\r
-    template int count_non_zero_caller<signed char>(const DevMem2D, PtrStep);\r
+    template int count_non_zero_caller<char>(const DevMem2D, PtrStep);\r
     template int count_non_zero_caller<unsigned short>(const DevMem2D, PtrStep);\r
-    template int count_non_zero_caller<signed short>(const DevMem2D, PtrStep);\r
+    template int count_non_zero_caller<short>(const DevMem2D, PtrStep);\r
     template int count_non_zero_caller<int>(const DevMem2D, PtrStep);\r
     template int count_non_zero_caller<float>(const DevMem2D, PtrStep);\r
     template int count_non_zero_caller<double>(const DevMem2D, PtrStep);\r
@@ -1171,9 +1171,9 @@ namespace cv { namespace gpu { namespace mathfunc
     }  \r
 \r
     template int count_non_zero_multipass_caller<unsigned char>(const DevMem2D, PtrStep);\r
-    template int count_non_zero_multipass_caller<signed char>(const DevMem2D, PtrStep);\r
+    template int count_non_zero_multipass_caller<char>(const DevMem2D, PtrStep);\r
     template int count_non_zero_multipass_caller<unsigned short>(const DevMem2D, PtrStep);\r
-    template int count_non_zero_multipass_caller<signed short>(const DevMem2D, PtrStep);\r
+    template int count_non_zero_multipass_caller<short>(const DevMem2D, PtrStep);\r
     template int count_non_zero_multipass_caller<int>(const DevMem2D, PtrStep);\r
     template int count_non_zero_multipass_caller<float>(const DevMem2D, PtrStep);\r
 \r

diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp
index bf212f6..2c94cb2 100644 (file)
--- a/tests/gpu/src/arithm.cpp
+++ b/tests/gpu/src/arithm.cpp
@@ -701,7 +701,7 @@ struct CV_GpuMinMaxTest: public CvTest
         for (int i = 0; i < src.rows; ++i)\r
         { \r
             Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i));\r
-            rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255));\r
+            rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256));\r
         }\r
 \r
         double minVal, maxVal;\r
@@ -714,7 +714,7 @@ struct CV_GpuMinMaxTest: public CvTest
         else \r
         {\r
             minVal = std::numeric_limits<double>::max();\r
-            maxVal = std::numeric_limits<double>::min();\r
+            maxVal = -std::numeric_limits<double>::max();\r
             for (int i = 0; i < src.rows; ++i)\r
                 for (int j = 0; j < src.cols; ++j)\r
                 {\r
@@ -747,7 +747,7 @@ struct CV_GpuMinMaxTest: public CvTest
         for (int i = 0; i < src.rows; ++i)\r
         { \r
             Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i));\r
-            rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255));\r
+            rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256));\r
         }\r
 \r
         cv::Mat mask(src.size(), CV_8U);\r
@@ -765,7 +765,7 @@ struct CV_GpuMinMaxTest: public CvTest
         {\r
             // OpenCV's minMaxLoc doesn't support CV_8S type \r
             minVal = std::numeric_limits<double>::max();\r
-            maxVal = std::numeric_limits<double>::min();\r
+            maxVal = -std::numeric_limits<double>::max();\r
             for (int i = 0; i < src_.rows; ++i)\r
                 for (int j = 0; j < src_.cols; ++j)\r
                 {\r
@@ -826,7 +826,7 @@ struct CV_GpuMinMaxLocTest: public CvTest
         for (int i = 0; i < src.rows; ++i)\r
         { \r
             Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i));\r
-            rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255));\r
+            rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256));\r
         }\r
 \r
         double minVal, maxVal;\r
@@ -838,7 +838,7 @@ struct CV_GpuMinMaxLocTest: public CvTest
         {\r
             // OpenCV's minMaxLoc doesn't support CV_8S type \r
             minVal = std::numeric_limits<double>::max();\r
-            maxVal = std::numeric_limits<double>::min();\r
+            maxVal = -std::numeric_limits<double>::max();\r
             for (int i = 0; i < src.rows; ++i)\r
                 for (int j = 0; j < src.cols; ++j)\r
                 {\r
@@ -895,7 +895,7 @@ struct CV_GpuCountNonZeroTest: CvTest
             for (int i = 0; i < src.rows; ++i)\r
             { \r
                 Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i));\r
-                rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255));\r
+                rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256));\r
             }\r
 \r
         int n_gold = cv::countNonZero(src);\r