diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index c7ca547976..48bed395d3 100644 --- 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 typedef void (*Caller)(const DevMem2D, double*, double*, PtrStep); static const Caller callers[2][7] = - { { min_max_multipass_caller, min_max_multipass_caller, - min_max_multipass_caller, min_max_multipass_caller, + { { min_max_multipass_caller, min_max_multipass_caller, + min_max_multipass_caller, min_max_multipass_caller, min_max_multipass_caller, min_max_multipass_caller, 0 }, - { min_max_caller, min_max_caller, - min_max_caller, min_max_caller, + { min_max_caller, min_max_caller, + min_max_caller, min_max_caller, min_max_caller, min_max_caller, min_max_caller } }; typedef void (*MaskedCaller)(const DevMem2D, const PtrStep, double*, double*, PtrStep); static const MaskedCaller masked_callers[2][7] = - { { min_max_mask_multipass_caller, min_max_mask_multipass_caller, - min_max_mask_multipass_caller, min_max_mask_multipass_caller, + { { min_max_mask_multipass_caller, min_max_mask_multipass_caller, + min_max_mask_multipass_caller, min_max_mask_multipass_caller, min_max_mask_multipass_caller, min_max_mask_multipass_caller, 0 }, - { min_max_mask_caller, min_max_mask_caller, - min_max_mask_caller, min_max_mask_caller, + { min_max_mask_caller, min_max_mask_caller, + min_max_mask_caller, min_max_mask_caller, min_max_mask_caller, min_max_mask_caller, min_max_mask_caller } }; @@ -615,9 +615,9 @@ void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point switch (src.type()) { case CV_8U: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; - case CV_8S: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; + case CV_8S: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; case CV_16U: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; - case CV_16S: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; + case CV_16S: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; case CV_32S: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; case CV_32F: min_max_loc_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; case CV_64F: @@ -634,9 +634,9 @@ void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point switch (src.type()) { case CV_8U: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; - case CV_8S: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; + case CV_8S: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; case CV_16U: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; - case CV_16S: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; + case CV_16S: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; case CV_32S: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; case CV_32F: min_max_loc_multipass_caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break; default: CV_Error(CV_StsBadArg, "minMaxLoc: unsupported type"); @@ -683,9 +683,9 @@ int cv::gpu::countNonZero(const GpuMat& src, GpuMat& buf) switch (src.type()) { case CV_8U: return count_non_zero_caller(src, buf); - case CV_8S: return count_non_zero_caller(src, buf); + case CV_8S: return count_non_zero_caller(src, buf); case CV_16U: return count_non_zero_caller(src, buf); - case CV_16S: return count_non_zero_caller(src, buf); + case CV_16S: return count_non_zero_caller(src, buf); case CV_32S: return count_non_zero_caller(src, buf); case CV_32F: return count_non_zero_caller(src, buf); case CV_64F: @@ -698,9 +698,9 @@ int cv::gpu::countNonZero(const GpuMat& src, GpuMat& buf) switch (src.type()) { case CV_8U: return count_non_zero_multipass_caller(src, buf); - case CV_8S: return count_non_zero_multipass_caller(src, buf); + case CV_8S: return count_non_zero_multipass_caller(src, buf); case CV_16U: return count_non_zero_multipass_caller(src, buf); - case CV_16S: return count_non_zero_multipass_caller(src, buf); + case CV_16S: return count_non_zero_multipass_caller(src, buf); case CV_32S: return count_non_zero_multipass_caller(src, buf); case CV_32F: return count_non_zero_multipass_caller(src, buf); } diff --git a/modules/gpu/src/cuda/limits_gpu.hpp b/modules/gpu/src/cuda/limits_gpu.hpp index 83a0743b55..fcf4bdc0b1 100644 --- 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(); } __device__ static type quiet_NaN() { return type(); } __device__ static type signaling_NaN() { return T(); } + static const bool is_signed; }; template<> struct numeric_limits_gpu @@ -71,6 +72,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = false; }; template<> struct numeric_limits_gpu @@ -84,6 +86,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = (char)-1 == -1; }; template<> struct numeric_limits_gpu @@ -97,6 +100,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = false; }; template<> struct numeric_limits_gpu @@ -110,6 +114,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = true; }; template<> struct numeric_limits_gpu @@ -123,6 +128,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = false; }; template<> struct numeric_limits_gpu @@ -136,6 +142,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = true; }; @@ -150,6 +157,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = false; }; template<> struct numeric_limits_gpu @@ -163,6 +171,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = true; }; template<> struct numeric_limits_gpu @@ -176,6 +185,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = false; }; template<> struct numeric_limits_gpu @@ -189,6 +199,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = true; }; template<> struct numeric_limits_gpu @@ -202,6 +213,7 @@ namespace cv __device__ static type infinity(); __device__ static type quiet_NaN(); __device__ static type signaling_NaN(); + static const bool is_signed = true; }; } } diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu index d4eceadd62..b00c94f546 100644 --- 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) template struct MinMaxTypeTraits {}; template <> struct MinMaxTypeTraits { typedef int best_type; }; - template <> struct MinMaxTypeTraits { typedef int best_type; }; + template <> struct MinMaxTypeTraits { typedef int best_type; }; template <> struct MinMaxTypeTraits { typedef int best_type; }; - template <> struct MinMaxTypeTraits { typedef int best_type; }; + template <> struct MinMaxTypeTraits { typedef int best_type; }; template <> struct MinMaxTypeTraits { typedef int best_type; }; template <> struct MinMaxTypeTraits { typedef float best_type; }; template <> struct MinMaxTypeTraits { typedef double best_type; }; @@ -492,7 +492,7 @@ namespace cv { namespace gpu { namespace mathfunc unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x; T mymin = numeric_limits_gpu::max(); - T mymax = numeric_limits_gpu::min(); + T mymax = numeric_limits_gpu::is_signed ? -numeric_limits_gpu::max() : numeric_limits_gpu::min(); unsigned int y_end = min(y0 + (ctheight - 1) * blockDim.y + 1, src.rows); unsigned int x_end = min(x0 + (ctwidth - 1) * blockDim.x + 1, src.cols); for (unsigned int y = y0; y < y_end; y += blockDim.y) @@ -584,9 +584,9 @@ namespace cv { namespace gpu { namespace mathfunc } template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); - template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); + template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); - template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); + template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); template void min_max_mask_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); @@ -613,9 +613,9 @@ namespace cv { namespace gpu { namespace mathfunc } template void min_max_caller(const DevMem2D, double*, double*, PtrStep); - template void min_max_caller(const DevMem2D, double*, double*, PtrStep); + template void min_max_caller(const DevMem2D, double*, double*, PtrStep); template void min_max_caller(const DevMem2D, double*, double*, PtrStep); - template void min_max_caller(const DevMem2D, double*, double*, PtrStep); + template void min_max_caller(const DevMem2D, double*, double*, PtrStep); template void min_max_caller(const DevMem2D, double*, double*, PtrStep); template void min_max_caller(const DevMem2D, double*,double*, PtrStep); template void min_max_caller(const DevMem2D, double*, double*, PtrStep); @@ -668,9 +668,9 @@ namespace cv { namespace gpu { namespace mathfunc } template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); - template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); + template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); - template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); + template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); template void min_max_mask_multipass_caller(const DevMem2D, const PtrStep, double*, double*, PtrStep); @@ -697,9 +697,9 @@ namespace cv { namespace gpu { namespace mathfunc } template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); - template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); + template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); - template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); + template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); template void min_max_multipass_caller(const DevMem2D, double*, double*, PtrStep); @@ -802,10 +802,10 @@ namespace cv { namespace gpu { namespace mathfunc unsigned int y0 = blockIdx.y * blockDim.y * ctheight + threadIdx.y; unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x; - T val = ((const T*)src.ptr(0))[0]; - T mymin = val, mymax = val; - unsigned int myminloc = 0, mymaxloc = 0; - + T mymin = numeric_limits_gpu::max(); + T mymax = numeric_limits_gpu::is_signed ? -numeric_limits_gpu::max() : numeric_limits_gpu::min(); + unsigned int myminloc = 0; + unsigned int mymaxloc = 0; unsigned int y_end = min(y0 + (ctheight - 1) * blockDim.y + 1, src.rows); unsigned int x_end = min(x0 + (ctwidth - 1) * blockDim.x + 1, src.cols); @@ -814,13 +814,13 @@ namespace cv { namespace gpu { namespace mathfunc const T* ptr = (const T*)src.ptr(y); for (unsigned int x = x0; x < x_end; x += blockDim.x) { - val = ptr[x]; - if (val < mymin) + T val = ptr[x]; + if (val <= mymin) { mymin = val; myminloc = y * src.cols + x; } - else if (val > mymax) + if (val >= mymax) { mymax = val; mymaxloc = y * src.cols + x; @@ -916,9 +916,9 @@ namespace cv { namespace gpu { namespace mathfunc } template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); - template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); + template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); - template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); + template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); template void min_max_loc_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); @@ -987,9 +987,9 @@ namespace cv { namespace gpu { namespace mathfunc } template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); - template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); + template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); - template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); + template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); template void min_max_loc_multipass_caller(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep); @@ -1126,9 +1126,9 @@ namespace cv { namespace gpu { namespace mathfunc } template int count_non_zero_caller(const DevMem2D, PtrStep); - template int count_non_zero_caller(const DevMem2D, PtrStep); + template int count_non_zero_caller(const DevMem2D, PtrStep); template int count_non_zero_caller(const DevMem2D, PtrStep); - template int count_non_zero_caller(const DevMem2D, PtrStep); + template int count_non_zero_caller(const DevMem2D, PtrStep); template int count_non_zero_caller(const DevMem2D, PtrStep); template int count_non_zero_caller(const DevMem2D, PtrStep); template int count_non_zero_caller(const DevMem2D, PtrStep); @@ -1171,9 +1171,9 @@ namespace cv { namespace gpu { namespace mathfunc } template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); - template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); + template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); - template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); + template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); template int count_non_zero_multipass_caller(const DevMem2D, PtrStep); diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp index bf212f645f..2c94cb277b 100644 --- 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) { Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i)); - rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255)); + rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256)); } double minVal, maxVal; @@ -714,7 +714,7 @@ struct CV_GpuMinMaxTest: public CvTest else { minVal = std::numeric_limits::max(); - maxVal = std::numeric_limits::min(); + maxVal = -std::numeric_limits::max(); for (int i = 0; i < src.rows; ++i) for (int j = 0; j < src.cols; ++j) { @@ -747,7 +747,7 @@ struct CV_GpuMinMaxTest: public CvTest for (int i = 0; i < src.rows; ++i) { Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i)); - rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255)); + rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256)); } cv::Mat mask(src.size(), CV_8U); @@ -765,7 +765,7 @@ struct CV_GpuMinMaxTest: public CvTest { // OpenCV's minMaxLoc doesn't support CV_8S type minVal = std::numeric_limits::max(); - maxVal = std::numeric_limits::min(); + maxVal = -std::numeric_limits::max(); for (int i = 0; i < src_.rows; ++i) for (int j = 0; j < src_.cols; ++j) { @@ -826,7 +826,7 @@ struct CV_GpuMinMaxLocTest: public CvTest for (int i = 0; i < src.rows; ++i) { Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i)); - rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255)); + rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256)); } double minVal, maxVal; @@ -838,7 +838,7 @@ struct CV_GpuMinMaxLocTest: public CvTest { // OpenCV's minMaxLoc doesn't support CV_8S type minVal = std::numeric_limits::max(); - maxVal = std::numeric_limits::min(); + maxVal = -std::numeric_limits::max(); for (int i = 0; i < src.rows; ++i) for (int j = 0; j < src.cols; ++j) { @@ -895,7 +895,7 @@ struct CV_GpuCountNonZeroTest: CvTest for (int i = 0; i < src.rows; ++i) { Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i)); - rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255)); + rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256)); } int n_gold = cv::countNonZero(src);