added masks support into gpu::minMaxLoc

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 13f237b..cdc9e4a 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -431,11 +431,12 @@ namespace cv
         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
+        CV_EXPORTS void minMaxLoc(const GpuMat& src, double* minVal, double* maxVal=0, Point* minLoc=0, Point* maxLoc=0, \r
+                                  const GpuMat& mask=GpuMat());\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, Point* minLoc, Point* maxLoc, \r
-                                  GpuMat& valbuf, GpuMat& locbuf);\r
+                                  const GpuMat& mask, GpuMat& valbuf, GpuMat& locbuf);\r
 \r
         //! counts non-zero array elements\r
         CV_EXPORTS int countNonZero(const GpuMat& src);\r

diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp
index 48bed39..54b01d7 100644 (file)
--- a/modules/gpu/src/arithm.cpp
+++ b/modules/gpu/src/arithm.cpp
@@ -67,8 +67,8 @@ void cv::gpu::flip(const GpuMat&, GpuMat&, int) { throw_nogpu(); }
 Scalar cv::gpu::sum(const GpuMat&) { throw_nogpu(); return Scalar(); }\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
+void cv::gpu::minMaxLoc(const GpuMat&, double*, double*, Point*, Point*, const GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::minMaxLoc(const GpuMat&, double*, double*, Point*, Point*, const GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); }\r
 int cv::gpu::countNonZero(const GpuMat&) { throw_nogpu(); return 0; }\r
 int cv::gpu::countNonZero(const GpuMat&, GpuMat&) { throw_nogpu(); return 0; }\r
 void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&) { throw_nogpu(); }\r
@@ -523,6 +523,8 @@ void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal, const Gp
     using namespace mathfunc::minmax;\r
 \r
     typedef void (*Caller)(const DevMem2D, double*, double*, PtrStep);\r
+    typedef void (*MaskedCaller)(const DevMem2D, const PtrStep, double*, double*, PtrStep);\r
+\r
     static const Caller callers[2][7] = \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
@@ -531,7 +533,6 @@ void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal, const Gp
         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<char>, \r
         min_max_mask_multipass_caller<unsigned short>, min_max_mask_multipass_caller<short>, \r
@@ -581,22 +582,53 @@ namespace cv { namespace gpu { namespace mathfunc { namespace minmaxloc {
                             int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf);\r
 \r
     template <typename T> \r
+    void min_max_loc_mask_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, \r
+                                 int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf);\r
+\r
+    template <typename T> \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
+                                     int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf);\r
+\r
+    template <typename T> \r
+    void min_max_loc_mask_multipass_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, \r
+                                           int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf);\r
+\r
 \r
 }}}}\r
 \r
-void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc)\r
+void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc, const GpuMat& mask)\r
 {    \r
     GpuMat valbuf, locbuf;\r
-    minMaxLoc(src, minVal, maxVal, minLoc, maxLoc, valbuf, locbuf);\r
+    minMaxLoc(src, minVal, maxVal, minLoc, maxLoc, mask, valbuf, locbuf);\r
 }\r
 \r
-void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc, \r
-                        GpuMat& valbuf, GpuMat& locbuf)\r
+void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc,\r
+                        const GpuMat& mask, GpuMat& valbuf, GpuMat& locbuf)\r
 {\r
     using namespace mathfunc::minmaxloc;\r
+\r
+    typedef void (*Caller)(const DevMem2D, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    typedef void (*MaskedCaller)(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+\r
+    static const Caller callers[2][7] = \r
+    { { min_max_loc_multipass_caller<unsigned char>, min_max_loc_multipass_caller<char>, \r
+        min_max_loc_multipass_caller<unsigned short>, min_max_loc_multipass_caller<short>, \r
+        min_max_loc_multipass_caller<int>, min_max_loc_multipass_caller<float>, 0 },\r
+      { min_max_loc_caller<unsigned char>, min_max_loc_caller<char>, \r
+        min_max_loc_caller<unsigned short>, min_max_loc_caller<short>, \r
+        min_max_loc_caller<int>, min_max_loc_caller<float>, min_max_loc_caller<double> } };\r
+\r
+    static const MaskedCaller masked_callers[2][7] = \r
+    { { min_max_loc_mask_multipass_caller<unsigned char>, min_max_loc_mask_multipass_caller<char>, \r
+        min_max_loc_mask_multipass_caller<unsigned short>, min_max_loc_mask_multipass_caller<short>, \r
+        min_max_loc_mask_multipass_caller<int>, min_max_loc_mask_multipass_caller<float>, 0 },\r
+      { min_max_loc_mask_caller<unsigned char>, min_max_loc_mask_caller<char>, \r
+        min_max_loc_mask_caller<unsigned short>, min_max_loc_mask_caller<short>, \r
+        min_max_loc_mask_caller<int>, min_max_loc_mask_caller<float>, min_max_loc_mask_caller<double> } };\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
@@ -609,38 +641,17 @@ void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point
     valbuf.create(valbuf_size, CV_8U);\r
     locbuf.create(locbuf_size, CV_8U);\r
 \r
-    int device = getDevice();\r
-    if (hasAtomicsSupport(device))\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<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<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
-            if (hasNativeDoubleSupport(device))\r
-            {\r
-                min_max_loc_caller<double>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); \r
-                break;\r
-            }\r
-        default: CV_Error(CV_StsBadArg, "minMaxLoc: unsupported type");\r
-        }\r
+    if (mask.empty())\r
+    {\r
+        Caller caller = callers[hasAtomicsSupport(getDevice())][src.type()];\r
+        if (!caller) CV_Error(CV_StsBadArg, "minMaxLoc: unsupported type");\r
+        caller(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf);\r
     }\r
     else\r
     {\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<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<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
+        MaskedCaller caller = masked_callers[hasAtomicsSupport(getDevice())][src.type()];\r
+        if (!caller) CV_Error(CV_StsBadArg, "minMaxLoc: unsupported type");\r
+        caller(src, mask, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf);\r
     }\r
 \r
     if (minLoc) { minLoc->x = minLoc_[0]; minLoc->y = minLoc_[1]; }\r
@@ -671,43 +682,27 @@ int cv::gpu::countNonZero(const GpuMat& src)
 int cv::gpu::countNonZero(const GpuMat& src, GpuMat& buf)\r
 {\r
     using namespace mathfunc::countnonzero;\r
+\r
+    typedef int (*Caller)(const DevMem2D src, PtrStep buf);\r
+\r
+    static const Caller callers[2][7] = \r
+    { { count_non_zero_multipass_caller<unsigned char>, count_non_zero_multipass_caller<char>,\r
+        count_non_zero_multipass_caller<unsigned short>, count_non_zero_multipass_caller<short>,\r
+        count_non_zero_multipass_caller<int>, count_non_zero_multipass_caller<float>, 0},\r
+      { count_non_zero_caller<unsigned char>, count_non_zero_caller<char>,\r
+        count_non_zero_caller<unsigned short>, count_non_zero_caller<short>,\r
+        count_non_zero_caller<int>, count_non_zero_caller<float>, count_non_zero_caller<double> } };\r
+\r
     CV_Assert(src.channels() == 1);\r
+    CV_Assert(src.type() != CV_64F || hasNativeDoubleSupport(getDevice()));\r
 \r
     Size buf_size;\r
     get_buf_size_required(buf_size.width, buf_size.height);\r
     buf.create(buf_size, CV_8U);\r
 \r
-    int device = getDevice();\r
-    if (hasAtomicsSupport(device))\r
-    {  \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<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<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
-            if (hasNativeDoubleSupport(device)) \r
-                return count_non_zero_caller<double>(src, buf);\r
-        }\r
-    }\r
-    else\r
-    {\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<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<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
-    CV_Error(CV_StsBadArg, "countNonZero: unsupported type");\r
-    return 0;\r
+    Caller caller = callers[hasAtomicsSupport(getDevice())][src.type()];\r
+    if (!caller) CV_Error(CV_StsBadArg, "countNonZero: unsupported type");\r
+    return caller(src, buf);\r
 }\r
 \r
 ////////////////////////////////////////////////////////////////////////\r

diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu
index b00c94f..fe52032 100644 (file)
--- a/modules/gpu/src/cuda/mathfunc.cu
+++ b/modules/gpu/src/cuda/mathfunc.cu
@@ -248,10 +248,10 @@ namespace cv { namespace gpu { namespace mathfunc
     struct Mask8U\r
     {\r
         explicit Mask8U(PtrStep mask): mask(mask) {}\r
-        __device__ bool operator()(int y, int x) { return mask.ptr(y)[x]; }\r
+        __device__ bool operator()(int y, int x) const { return mask.ptr(y)[x]; }\r
         PtrStep mask;\r
     };\r
-    struct MaskTrue { __device__ bool operator()(int y, int x) { return true; } };\r
+    struct MaskTrue { __device__ bool operator()(int y, int x) const { return true; } };\r
 \r
     // Unary operations\r
 \r
@@ -788,8 +788,8 @@ namespace cv { namespace gpu { namespace mathfunc
     }\r
 \r
 \r
-    template <int nthreads, typename T>\r
-    __global__ void min_max_loc_kernel(const DevMem2D src, T* minval, T* maxval, \r
+    template <int nthreads, typename T, typename Mask>\r
+    __global__ void min_max_loc_kernel(const DevMem2D src, Mask mask, T* minval, T* maxval, \r
                                        unsigned int* minloc, unsigned int* maxloc)\r
     {\r
         typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
@@ -814,16 +814,11 @@ 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
-                T val = ptr[x];\r
-                if (val <= mymin) \r
-                { \r
-                    mymin = val; \r
-                    myminloc = y * src.cols + x; \r
-                }\r
-                if (val >= mymax)\r
+                if (mask(y, x))\r
                 {\r
-                    mymax = val; \r
-                    mymaxloc = y * src.cols + x; \r
+                    T val = ptr[x];\r
+                    if (val <= mymin) { mymin = val; myminloc = y * src.cols + x; }\r
+                    if (val >= mymax) { mymax = val; mymaxloc = y * src.cols + x; }\r
                 }\r
             }\r
         }\r
@@ -887,6 +882,44 @@ namespace cv { namespace gpu { namespace mathfunc
 \r
 \r
     template <typename T>\r
+    void min_max_loc_mask_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, \r
+                                 int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf)\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*)valbuf.ptr(0);\r
+        T* maxval_buf = (T*)valbuf.ptr(1);\r
+        unsigned int* minloc_buf = (unsigned int*)locbuf.ptr(0);\r
+        unsigned int* maxloc_buf = (unsigned int*)locbuf.ptr(1);\r
+\r
+        min_max_loc_kernel<256, T, Mask8U><<<grid, threads>>>(src, Mask8U(mask), minval_buf, maxval_buf, minloc_buf, maxloc_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
+        unsigned int minloc_, maxloc_;\r
+        cudaSafeCall(cudaMemcpy(&minloc_, minloc_buf, sizeof(int), cudaMemcpyDeviceToHost));\r
+        cudaSafeCall(cudaMemcpy(&maxloc_, maxloc_buf, sizeof(int), cudaMemcpyDeviceToHost));\r
+        minloc[1] = minloc_ / src.cols; minloc[0] = minloc_ - minloc[1] * src.cols;\r
+        maxloc[1] = maxloc_ / src.cols; maxloc[0] = maxloc_ - maxloc[1] * src.cols;\r
+    }\r
+\r
+    template void min_max_loc_mask_caller<unsigned char>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_caller<char>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_caller<unsigned short>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_caller<short>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_caller<int>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_caller<float>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_caller<double>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+\r
+\r
+    template <typename T>\r
     void min_max_loc_caller(const DevMem2D src, double* minval, double* maxval, \r
                             int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf)\r
     {\r
@@ -899,7 +932,7 @@ namespace cv { namespace gpu { namespace mathfunc
         unsigned int* minloc_buf = (unsigned int*)locbuf.ptr(0);\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<256, T, MaskTrue><<<grid, threads>>>(src, MaskTrue(), minval_buf, maxval_buf, minloc_buf, maxloc_buf);\r
         cudaSafeCall(cudaThreadSynchronize());\r
 \r
         T minval_, maxval_;\r
@@ -957,8 +990,46 @@ namespace cv { namespace gpu { namespace mathfunc
 \r
 \r
     template <typename T>\r
+    void min_max_loc_mask_multipass_caller(const DevMem2D src, const PtrStep mask, double* minval, double* maxval, \r
+                                           int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf)\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*)valbuf.ptr(0);\r
+        T* maxval_buf = (T*)valbuf.ptr(1);\r
+        unsigned int* minloc_buf = (unsigned int*)locbuf.ptr(0);\r
+        unsigned int* maxloc_buf = (unsigned int*)locbuf.ptr(1);\r
+\r
+        min_max_loc_kernel<256, T, Mask8U><<<grid, threads>>>(src, Mask8U(mask), minval_buf, maxval_buf, minloc_buf, maxloc_buf);\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
+        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
+        unsigned int minloc_, maxloc_;\r
+        cudaSafeCall(cudaMemcpy(&minloc_, minloc_buf, sizeof(int), cudaMemcpyDeviceToHost));\r
+        cudaSafeCall(cudaMemcpy(&maxloc_, maxloc_buf, sizeof(int), cudaMemcpyDeviceToHost));\r
+        minloc[1] = minloc_ / src.cols; minloc[0] = minloc_ - minloc[1] * src.cols;\r
+        maxloc[1] = maxloc_ / src.cols; maxloc[0] = maxloc_ - maxloc[1] * src.cols;\r
+    }\r
+\r
+    template void min_max_loc_mask_multipass_caller<unsigned char>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_multipass_caller<char>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_multipass_caller<unsigned short>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_multipass_caller<short>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_multipass_caller<int>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+    template void min_max_loc_mask_multipass_caller<float>(const DevMem2D, const PtrStep, double*, double*, int[2], int[2], PtrStep, PtrStep);\r
+\r
+\r
+    template <typename T>\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
+                                      int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf)\r
     {\r
         dim3 threads, grid;\r
         estimate_thread_cfg(threads, grid);\r
@@ -969,7 +1040,7 @@ namespace cv { namespace gpu { namespace mathfunc
         unsigned int* minloc_buf = (unsigned int*)locbuf.ptr(0);\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<256, T, MaskTrue><<<grid, threads>>>(src, MaskTrue(), minval_buf, maxval_buf, minloc_buf, maxloc_buf);\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

diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp
index 2c94cb2..02ff33b 100644 (file)
--- a/tests/gpu/src/arithm.cpp
+++ b/tests/gpu/src/arithm.cpp
@@ -684,7 +684,7 @@ struct CV_GpuMinMaxTest: public CvTest
         if (cv::gpu::hasNativeDoubleSupport(cv::gpu::getDevice())) depth_end = CV_64F; else depth_end = CV_32F;\r
         for (int depth = CV_8U; depth <= depth_end; ++depth)\r
         {\r
-            for (int i = 0; i < 1; ++i)\r
+            for (int i = 0; i < 3; ++i)\r
             {\r
                 int rows = 1 + rand() % 1000;\r
                 int cols = 1 + rand() % 1000;\r
@@ -829,11 +829,14 @@ struct CV_GpuMinMaxLocTest: public CvTest
             rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(256));\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
         if (depth != CV_8S)       \r
-            cv::minMaxLoc(src, &minVal, &maxVal, &minLoc, &maxLoc);\r
+            cv::minMaxLoc(src, &minVal, &maxVal, &minLoc, &maxLoc, mask);\r
         else \r
         {\r
             // OpenCV's minMaxLoc doesn't support CV_8S type \r
@@ -843,14 +846,17 @@ struct CV_GpuMinMaxLocTest: public CvTest
                 for (int j = 0; j < src.cols; ++j)\r
                 {\r
                     char val = src.at<char>(i, j);\r
-                    if (val < minVal) { minVal = val; minLoc = cv::Point(j, i); }\r
-                    if (val > maxVal) { maxVal = val; maxLoc = cv::Point(j, i); }\r
+                    if (mask.at<unsigned char>(i, j))\r
+                    {\r
+                        if (val < minVal) { minVal = val; minLoc = cv::Point(j, i); }\r
+                        if (val > maxVal) { maxVal = val; maxLoc = cv::Point(j, i); }\r
+                    }\r
                 }\r
         }\r
 \r
         double minVal_, maxVal_;\r
         cv::Point minLoc_, maxLoc_;        \r
-        cv::gpu::minMaxLoc(cv::gpu::GpuMat(src), &minVal_, &maxVal_, &minLoc_, &maxLoc_, valbuf, locbuf);\r
+        cv::gpu::minMaxLoc(cv::gpu::GpuMat(src), &minVal_, &maxVal_, &minLoc_, &maxLoc_, cv::gpu::GpuMat(mask), valbuf, locbuf);\r
        \r
         CHECK(minVal == minVal_, CvTS::FAIL_INVALID_OUTPUT);\r
         CHECK(maxVal == maxVal_, CvTS::FAIL_INVALID_OUTPUT);\r