\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
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
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
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
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
__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
__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
__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
__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
__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
__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
__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
__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
__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
__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
__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
__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
// 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
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
} \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
} \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
}\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
}\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
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
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
}\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
}\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
} \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
} \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
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
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
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
{\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
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
{\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
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