CV_EXPORTS Scalar sum(const GpuMat& m);\r
\r
//! finds global minimum and maximum array elements and returns their values\r
- CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal=0);\r
+ CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal=0, const GpuMat& mask=GpuMat());\r
\r
//! finds global minimum and maximum array elements and returns their values\r
- CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal, GpuMat& buf);\r
+ CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal, const GpuMat& mask, GpuMat& buf);\r
\r
//! finds global minimum and maximum array elements and returns their values with locations\r
CV_EXPORTS void minMaxLoc(const GpuMat& src, double* minVal, double* maxVal=0, Point* minLoc=0, Point* maxLoc=0);\r
double cv::gpu::norm(const GpuMat&, const GpuMat&, int) { throw_nogpu(); return 0.0; }\r
void cv::gpu::flip(const GpuMat&, GpuMat&, int) { throw_nogpu(); }\r
Scalar cv::gpu::sum(const GpuMat&) { throw_nogpu(); return Scalar(); }\r
-void cv::gpu::minMax(const GpuMat&, double*, double*) { throw_nogpu(); }\r
-void cv::gpu::minMax(const GpuMat&, double*, double*, GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::minMax(const GpuMat&, double*, double*, const GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::minMax(const GpuMat&, double*, double*, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
void cv::gpu::minMaxLoc(const GpuMat&, double*, double*, Point*, Point*) { throw_nogpu(); }\r
void cv::gpu::minMaxLoc(const GpuMat&, double*, double*, Point*, Point*, GpuMat&, GpuMat&) { throw_nogpu(); }\r
int cv::gpu::countNonZero(const GpuMat&) { throw_nogpu(); return 0; }\r
void min_max_caller(const DevMem2D src, double* minval, double* maxval, PtrStep buf);\r
\r
template <typename T> \r
- void min_max_caller_2steps(const DevMem2D src, double* minval, double* maxval, PtrStep buf);\r
+ void min_max_mask_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, PtrStep buf);\r
+\r
+ template <typename T> \r
+ void min_max_multipass_caller(const DevMem2D src, double* minval, double* maxval, PtrStep buf);\r
+\r
+ template <typename T> \r
+ void min_max_mask_multipass_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, PtrStep buf);\r
\r
}}}}\r
\r
-void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal)\r
+void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal, const GpuMat& mask)\r
{\r
GpuMat buf;\r
- minMax(src, minVal, maxVal, buf);\r
+ minMax(src, minVal, maxVal, mask, buf);\r
}\r
\r
-\r
-void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal, GpuMat& buf)\r
+void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal, const GpuMat& mask, GpuMat& buf)\r
{\r
using namespace mathfunc::minmax;\r
\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<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<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<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<int>, min_max_mask_caller<float>, \r
+ min_max_mask_caller<double> } };\r
+\r
+\r
+ CV_Assert(src.channels() == 1);\r
+ CV_Assert(mask.empty() || (mask.type() == CV_8U && src.size() == mask.size()));\r
+ CV_Assert(src.type() != CV_64F || hasNativeDoubleSupport(getDevice()));\r
+\r
double minVal_; if (!minVal) minVal = &minVal_;\r
double maxVal_; if (!maxVal) maxVal = &maxVal_;\r
-\r
- GpuMat src_ = src.reshape(1);\r
\r
Size bufSize;\r
get_buf_size_required(src.elemSize(), bufSize.width, bufSize.height);\r
buf.create(bufSize, CV_8U);\r
\r
- int device = getDevice();\r
- if (hasAtomicsSupport(device))\r
+ if (mask.empty())\r
{\r
- switch (src_.type())\r
- {\r
- case CV_8U: min_max_caller<unsigned char>(src_, minVal, maxVal, buf); break;\r
- case CV_8S: min_max_caller<signed char>(src_, minVal, maxVal, buf); break;\r
- case CV_16U: min_max_caller<unsigned short>(src_, minVal, maxVal, buf); break;\r
- case CV_16S: min_max_caller<signed short>(src_, minVal, maxVal, buf); break;\r
- case CV_32S: min_max_caller<int>(src_, minVal, maxVal, buf); break;\r
- case CV_32F: min_max_caller<float>(src_, minVal, maxVal, buf); break;\r
- case CV_64F: \r
- if (hasNativeDoubleSupport(device))\r
- {\r
- min_max_caller<double>(src_, minVal, maxVal, buf); \r
- break;\r
- }\r
- default: CV_Error(CV_StsBadArg, "minMax: unsupported type");\r
- }\r
+ Caller caller = callers[hasAtomicsSupport(getDevice())][src.type()];\r
+ if (!caller) CV_Error(CV_StsBadArg, "minMax: unsupported type");\r
+ caller(src, minVal, maxVal, buf);\r
}\r
else\r
{\r
- switch (src_.type())\r
- {\r
- case CV_8U: min_max_caller_2steps<unsigned char>(src_, minVal, maxVal, buf); break;\r
- case CV_8S: min_max_caller_2steps<signed char>(src_, minVal, maxVal, buf); break;\r
- case CV_16U: min_max_caller_2steps<unsigned short>(src_, minVal, maxVal, buf); break;\r
- case CV_16S: min_max_caller_2steps<signed short>(src_, minVal, maxVal, buf); break;\r
- case CV_32S: min_max_caller_2steps<int>(src_, minVal, maxVal, buf); break;\r
- case CV_32F: min_max_caller_2steps<float>(src_, minVal, maxVal, buf); break;\r
- default: CV_Error(CV_StsBadArg, "minMax: unsupported type");\r
- }\r
+ MaskedCaller caller = masked_callers[hasAtomicsSupport(getDevice())][src.type()];\r
+ if (!caller) CV_Error(CV_StsBadArg, "minMax: unsupported type");\r
+ caller(src, mask, minVal, maxVal, buf);\r
}\r
}\r
\r
int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf);\r
\r
template <typename T> \r
- void min_max_loc_caller_2steps(const DevMem2D src, double* minval, double* maxval, \r
+ void min_max_loc_multipass_caller(const DevMem2D src, double* minval, double* maxval, \r
int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf);\r
\r
}}}}\r
{\r
switch (src.type())\r
{\r
- case CV_8U: min_max_loc_caller_2steps<unsigned char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
- case CV_8S: min_max_loc_caller_2steps<signed char>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
- case CV_16U: min_max_loc_caller_2steps<unsigned short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
- case CV_16S: min_max_loc_caller_2steps<signed short>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
- case CV_32S: min_max_loc_caller_2steps<int>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
- case CV_32F: min_max_loc_caller_2steps<float>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\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_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_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
}\r
}\r
int count_non_zero_caller(const DevMem2D src, PtrStep buf);\r
\r
template <typename T> \r
- int count_non_zero_caller_2steps(const DevMem2D src, PtrStep buf);\r
+ int count_non_zero_multipass_caller(const DevMem2D src, PtrStep buf);\r
\r
}}}}\r
\r
{\r
switch (src.type())\r
{\r
- case CV_8U: return count_non_zero_caller_2steps<unsigned char>(src, buf);\r
- case CV_8S: return count_non_zero_caller_2steps<signed char>(src, buf);\r
- case CV_16U: return count_non_zero_caller_2steps<unsigned short>(src, buf);\r
- case CV_16S: return count_non_zero_caller_2steps<signed short>(src, buf);\r
- case CV_32S: return count_non_zero_caller_2steps<int>(src, buf);\r
- case CV_32F: return count_non_zero_caller_2steps<float>(src, buf);\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_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_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
}\r
\r
}\r
\r
\r
- template <int nthreads, typename T>\r
- __global__ void min_max_kernel(const DevMem2D src, T* minval, T* maxval)\r
+ template <int nthreads, typename T, typename Mask>\r
+ __global__ void min_max_kernel(const DevMem2D src, Mask mask, T* minval, T* maxval)\r
{\r
typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
__shared__ best_type sminval[nthreads];\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;\r
T mymin = numeric_limits_gpu<T>::max();\r
T mymax = numeric_limits_gpu<T>::min();\r
- for (unsigned int y = 0; y < ctheight && y0 + y * blockDim.y < src.rows; ++y)\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
- const T* ptr = (const T*)src.ptr(y0 + y * blockDim.y);\r
- for (unsigned int x = 0; x < ctwidth && x0 + x * blockDim.x < src.cols; ++x)\r
+ const T* src_row = (const T*)src.ptr(y);\r
+ for (unsigned int x = x0; x < x_end; x += blockDim.x)\r
{\r
- val = ptr[x0 + x * blockDim.x];\r
- mymin = min(mymin, val);\r
- mymax = max(mymax, val);\r
+ T val = src_row[x];\r
+ if (mask(y, x)) \r
+ { \r
+ mymin = min(mymin, val); \r
+ mymax = max(mymax, val); \r
+ }\r
}\r
}\r
\r
\r
\r
template <typename T>\r
+ void min_max_mask_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, PtrStep buf)\r
+ {\r
+ dim3 threads, grid;\r
+ estimate_thread_cfg(threads, grid);\r
+ estimate_kernel_consts(src.cols, src.rows, threads, grid);\r
+\r
+ T* minval_buf = (T*)buf.ptr(0);\r
+ T* maxval_buf = (T*)buf.ptr(1);\r
+\r
+ min_max_kernel<256, T, Mask8U><<<grid, threads>>>(src, Mask8U(mask), minval_buf, maxval_buf);\r
+ cudaSafeCall(cudaThreadSynchronize());\r
+\r
+ T minval_, maxval_;\r
+ cudaSafeCall(cudaMemcpy(&minval_, minval_buf, sizeof(T), cudaMemcpyDeviceToHost));\r
+ cudaSafeCall(cudaMemcpy(&maxval_, maxval_buf, sizeof(T), cudaMemcpyDeviceToHost));\r
+ *minval = minval_;\r
+ *maxval = maxval_;\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<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<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 <typename T>\r
void min_max_caller(const DevMem2D src, double* minval, double* maxval, PtrStep buf)\r
{\r
dim3 threads, grid;\r
T* minval_buf = (T*)buf.ptr(0);\r
T* maxval_buf = (T*)buf.ptr(1);\r
\r
- min_max_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf);\r
+ min_max_kernel<256, T, MaskTrue><<<grid, threads>>>(src, MaskTrue(), minval_buf, maxval_buf);\r
cudaSafeCall(cudaThreadSynchronize());\r
\r
T minval_, maxval_;\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<int>(const DevMem2D, double*, double*, PtrStep);\r
- template void min_max_caller<float>(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
- // This kernel will be used only when compute capability is 1.0\r
template <int nthreads, typename T>\r
- __global__ void min_max_kernel_2ndstep(T* minval, T* maxval, int size)\r
+ __global__ void min_max_pass2_kernel(T* minval, T* maxval, int size)\r
{\r
typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
__shared__ best_type sminval[nthreads];\r
\r
\r
template <typename T>\r
- void min_max_caller_2steps(const DevMem2D src, double* minval, double* maxval, PtrStep buf)\r
+ void min_max_mask_multipass_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, PtrStep buf)\r
+ {\r
+ dim3 threads, grid;\r
+ estimate_thread_cfg(threads, grid);\r
+ estimate_kernel_consts(src.cols, src.rows, threads, grid);\r
+\r
+ T* minval_buf = (T*)buf.ptr(0);\r
+ T* maxval_buf = (T*)buf.ptr(1);\r
+\r
+ min_max_kernel<256, T, Mask8U><<<grid, threads>>>(src, Mask8U(mask), minval_buf, maxval_buf);\r
+ min_max_pass2_kernel<256, T><<<1, 256>>>(minval_buf, maxval_buf, grid.x * grid.y);\r
+ cudaSafeCall(cudaThreadSynchronize());\r
+\r
+ T minval_, maxval_;\r
+ cudaSafeCall(cudaMemcpy(&minval_, minval_buf, sizeof(T), cudaMemcpyDeviceToHost));\r
+ cudaSafeCall(cudaMemcpy(&maxval_, maxval_buf, sizeof(T), cudaMemcpyDeviceToHost));\r
+ *minval = minval_;\r
+ *maxval = maxval_;\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<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<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
+ template <typename T>\r
+ void min_max_multipass_caller(const DevMem2D src, double* minval, double* maxval, PtrStep buf)\r
{\r
dim3 threads, grid;\r
estimate_thread_cfg(threads, grid);\r
T* minval_buf = (T*)buf.ptr(0);\r
T* maxval_buf = (T*)buf.ptr(1);\r
\r
- min_max_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf);\r
- min_max_kernel_2ndstep<256, T><<<1, 256>>>(minval_buf, maxval_buf, grid.x * grid.y);\r
+ min_max_kernel<256, T, MaskTrue><<<grid, threads>>>(src, MaskTrue(), minval_buf, maxval_buf);\r
+ min_max_pass2_kernel<256, T><<<1, 256>>>(minval_buf, maxval_buf, grid.x * grid.y);\r
cudaSafeCall(cudaThreadSynchronize());\r
\r
T minval_, maxval_;\r
*maxval = maxval_;\r
}\r
\r
- template void min_max_caller_2steps<unsigned char>(const DevMem2D, double*, double*, PtrStep);\r
- template void min_max_caller_2steps<signed char>(const DevMem2D, double*, double*, PtrStep);\r
- template void min_max_caller_2steps<unsigned short>(const DevMem2D, double*, double*, PtrStep);\r
- template void min_max_caller_2steps<signed short>(const DevMem2D, double*, double*, PtrStep);\r
- template void min_max_caller_2steps<int>(const DevMem2D, double*, double*, PtrStep);\r
- template void min_max_caller_2steps<float>(const DevMem2D, double*, double*, PtrStep);\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<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<int>(const DevMem2D, double*, double*, PtrStep);\r
+ template void min_max_multipass_caller<float>(const DevMem2D, double*, double*, PtrStep);\r
\r
} // namespace minmax\r
\r
\r
// This kernel will be used only when compute capability is 1.0\r
template <int nthreads, typename T>\r
- __global__ void min_max_loc_kernel_2ndstep(T* minval, T* maxval, unsigned int* minloc, unsigned int* maxloc, int size)\r
+ __global__ void min_max_loc_pass2_kernel(T* minval, T* maxval, unsigned int* minloc, unsigned int* maxloc, int size)\r
{\r
typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
__shared__ best_type sminval[nthreads];\r
\r
\r
template <typename T>\r
- void min_max_loc_caller_2steps(const DevMem2D src, double* minval, double* maxval, \r
+ void min_max_loc_multipass_caller(const DevMem2D src, double* minval, double* maxval, \r
int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf)\r
{\r
dim3 threads, grid;\r
unsigned int* maxloc_buf = (unsigned int*)locbuf.ptr(1);\r
\r
min_max_loc_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf, minloc_buf, maxloc_buf);\r
- min_max_loc_kernel_2ndstep<256, T><<<1, 256>>>(minval_buf, maxval_buf, minloc_buf, maxloc_buf, grid.x * grid.y);\r
+ min_max_loc_pass2_kernel<256, T><<<1, 256>>>(minval_buf, maxval_buf, minloc_buf, maxloc_buf, grid.x * grid.y);\r
cudaSafeCall(cudaThreadSynchronize());\r
\r
T minval_, maxval_;\r
maxloc[1] = maxloc_ / src.cols; maxloc[0] = maxloc_ - maxloc[1] * src.cols;\r
}\r
\r
- template void min_max_loc_caller_2steps<unsigned char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
- template void min_max_loc_caller_2steps<signed char>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
- template void min_max_loc_caller_2steps<unsigned short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
- template void min_max_loc_caller_2steps<signed short>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
- template void min_max_loc_caller_2steps<int>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
- template void min_max_loc_caller_2steps<float>(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\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<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<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
} // namespace minmaxloc\r
\r
\r
\r
template <int nthreads, typename T>\r
- __global__ void count_non_zero_kernel_2ndstep(unsigned int* count, int size)\r
+ __global__ void count_non_zero_pass2_kernel(unsigned int* count, int size)\r
{\r
__shared__ unsigned int scount[nthreads];\r
unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
\r
\r
template <typename T>\r
- int count_non_zero_caller_2steps(const DevMem2D src, PtrStep buf)\r
+ int count_non_zero_multipass_caller(const DevMem2D src, PtrStep buf)\r
{\r
dim3 threads, grid;\r
estimate_thread_cfg(threads, grid);\r
unsigned int* count_buf = (unsigned int*)buf.ptr(0);\r
\r
count_non_zero_kernel<256, T><<<grid, threads>>>(src, count_buf);\r
- count_non_zero_kernel_2ndstep<256, T><<<1, 256>>>(count_buf, grid.x * grid.y);\r
+ count_non_zero_pass2_kernel<256, T><<<1, 256>>>(count_buf, grid.x * grid.y);\r
cudaSafeCall(cudaThreadSynchronize());\r
\r
unsigned int count;\r
return count;\r
} \r
\r
- template int count_non_zero_caller_2steps<unsigned char>(const DevMem2D, PtrStep);\r
- template int count_non_zero_caller_2steps<signed char>(const DevMem2D, PtrStep);\r
- template int count_non_zero_caller_2steps<unsigned short>(const DevMem2D, PtrStep);\r
- template int count_non_zero_caller_2steps<signed short>(const DevMem2D, PtrStep);\r
- template int count_non_zero_caller_2steps<int>(const DevMem2D, PtrStep);\r
- template int count_non_zero_caller_2steps<float>(const DevMem2D, PtrStep);\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<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<int>(const DevMem2D, PtrStep);\r
+ template int count_non_zero_multipass_caller<float>(const DevMem2D, PtrStep);\r
\r
} // namespace countnonzero\r
\r
{\r
int depth_end;\r
if (cv::gpu::hasNativeDoubleSupport(cv::gpu::getDevice())) depth_end = CV_64F; else depth_end = CV_32F;\r
- for (int cn = 1; cn <= 4; ++cn)\r
- for (int depth = CV_8U; depth <= depth_end; ++depth)\r
+ for (int depth = CV_8U; depth <= depth_end; ++depth)\r
+ {\r
+ for (int i = 0; i < 1; ++i)\r
{\r
- for (int i = 0; i < 1; ++i)\r
- {\r
- int rows = 1 + rand() % 1000;\r
- int cols = 1 + rand() % 1000;\r
- test(rows, cols, cn, depth);\r
- }\r
+ int rows = 1 + rand() % 1000;\r
+ int cols = 1 + rand() % 1000;\r
+ test(rows, cols, 1, depth);\r
+ test_masked(rows, cols, 1, depth);\r
}\r
+ }\r
}\r
\r
void test(int rows, int cols, int cn, int depth)\r
double minVal, maxVal;\r
cv::Point minLoc, maxLoc;\r
\r
+ if (depth != CV_8S)\r
+ {\r
+ cv::minMaxLoc(src, &minVal, &maxVal, &minLoc, &maxLoc);\r
+ }\r
+ else \r
+ {\r
+ minVal = std::numeric_limits<double>::max();\r
+ maxVal = std::numeric_limits<double>::min();\r
+ for (int i = 0; i < src.rows; ++i)\r
+ for (int j = 0; j < src.cols; ++j)\r
+ {\r
+ signed char val = src.at<signed char>(i, j);\r
+ if (val < minVal) minVal = val;\r
+ if (val > maxVal) maxVal = val;\r
+ }\r
+ }\r
+\r
+ double minVal_, maxVal_;\r
+ cv::Point minLoc_, maxLoc_; \r
+ cv::gpu::minMax(cv::gpu::GpuMat(src), &minVal_, &maxVal_, cv::gpu::GpuMat(), buf);\r
+ \r
+ if (abs(minVal - minVal_) > 1e-3f)\r
+ {\r
+ ts->printf(CvTS::CONSOLE, "\nfail: minVal=%f minVal_=%f rows=%d cols=%d depth=%d cn=%d\n", minVal, minVal_, rows, cols, depth, cn);\r
+ ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);\r
+ }\r
+ if (abs(maxVal - maxVal_) > 1e-3f)\r
+ {\r
+ ts->printf(CvTS::CONSOLE, "\nfail: maxVal=%f maxVal_=%f rows=%d cols=%d depth=%d cn=%d\n", maxVal, maxVal_, rows, cols, depth, cn);\r
+ ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);\r
+ }\r
+ } \r
+\r
+ void test_masked(int rows, int cols, int cn, int depth)\r
+ {\r
+ cv::Mat src(rows, cols, CV_MAKE_TYPE(depth, cn));\r
+ cv::RNG rng;\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
+ }\r
+\r
+ cv::Mat mask(src.size(), CV_8U);\r
+ rng.fill(mask, RNG::UNIFORM, Scalar(0), Scalar(2));\r
+\r
+ double minVal, maxVal;\r
+ cv::Point minLoc, maxLoc;\r
+\r
Mat src_ = src.reshape(1);\r
if (depth != CV_8S)\r
{\r
- cv::minMaxLoc(src_, &minVal, &maxVal, &minLoc, &maxLoc);\r
+ cv::minMaxLoc(src_, &minVal, &maxVal, &minLoc, &maxLoc, mask);\r
}\r
else \r
{\r
for (int j = 0; j < src_.cols; ++j)\r
{\r
char val = src_.at<char>(i, j);\r
- if (val < minVal) minVal = val;\r
- if (val > maxVal) maxVal = val;\r
+ if (mask.at<unsigned char>(i, j)) { if (val < minVal) minVal = val; }\r
+ if (mask.at<unsigned char>(i, j)) { if (val > maxVal) maxVal = val; }\r
}\r
}\r
\r
double minVal_, maxVal_;\r
cv::Point minLoc_, maxLoc_; \r
- cv::gpu::minMax(cv::gpu::GpuMat(src), &minVal_, &maxVal_, buf);\r
+ cv::gpu::minMax(cv::gpu::GpuMat(src), &minVal_, &maxVal_, cv::gpu::GpuMat(mask), buf);\r
\r
if (abs(minVal - minVal_) > 1e-3f)\r
{\r