added minMaxLoc function into gpu module

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 70f9eda..4dcbd0b 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -422,7 +422,10 @@ namespace cv
         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);\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
 \r
         //! transforms 8-bit unsigned integers using lookup table: dst(i)=lut(src(i))\r
         //! destination array will have the depth type as lut and the same channels number as source\r

diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp
index 7c54719..d74fbb2 100644 (file)
--- a/modules/gpu/src/arithm.cpp
+++ b/modules/gpu/src/arithm.cpp
@@ -66,6 +66,7 @@ double cv::gpu::norm(const GpuMat&, const GpuMat&, int) { throw_nogpu(); return
 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::minMaxLoc(const GpuMat&, double*, double*, Point*, Point*) { throw_nogpu(); }\r
 void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&) { throw_nogpu(); }\r
 void cv::gpu::exp(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
 void cv::gpu::log(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
@@ -530,6 +531,57 @@ void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal)
     }\r
 }\r
 \r
+\r
+////////////////////////////////////////////////////////////////////////\r
+// minMaxLoc\r
+\r
+namespace cv { namespace gpu { namespace mathfunc {\r
+    template <typename T> \r
+    void min_max_loc_caller(const DevMem2D src, double* minval, double* maxval, int* minlocx, int* minlocy,\r
+                                                                                int* maxlocx, int* maxlocy);\r
+}}}\r
+\r
+void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc)\r
+{\r
+    CV_Assert(src.channels() == 1);\r
+\r
+    double maxVal_;\r
+    if (!maxVal) maxVal = &maxVal_;\r
+\r
+    cv::Point minLoc_;\r
+    if (!minLoc) minLoc = &minLoc_;\r
+\r
+    cv::Point maxLoc_;\r
+    if (!maxLoc) maxLoc = &maxLoc_;\r
+  \r
+    switch (src.type())\r
+    {\r
+    case CV_8U:\r
+        mathfunc::min_max_loc_caller<unsigned char>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
+        break;\r
+    case CV_8S:\r
+        mathfunc::min_max_loc_caller<signed char>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
+        break;\r
+    case CV_16U:\r
+        mathfunc::min_max_loc_caller<unsigned short>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
+        break;\r
+    case CV_16S:\r
+        mathfunc::min_max_loc_caller<signed short>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
+        break;\r
+    case CV_32S:\r
+        mathfunc::min_max_loc_caller<int>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
+        break;\r
+    case CV_32F:\r
+        mathfunc::min_max_loc_caller<float>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
+        break;\r
+    case CV_64F:\r
+        mathfunc::min_max_loc_caller<double>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
+        break;\r
+    default:\r
+        CV_Error(CV_StsBadArg, "Unsupported type");\r
+    }\r
+}\r
+\r
 ////////////////////////////////////////////////////////////////////////\r
 // LUT\r
 \r

diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu
index d69a32a..bff3a30 100644 (file)
--- a/modules/gpu/src/cuda/mathfunc.cu
+++ b/modules/gpu/src/cuda/mathfunc.cu
@@ -410,10 +410,10 @@ namespace cv { namespace gpu { namespace mathfunc
     template <> struct MinMaxTypeTraits<float> { typedef float best_type; };\r
     template <> struct MinMaxTypeTraits<double> { typedef double best_type; };\r
 \r
-    template <typename T, int op> struct Cmp {};\r
+    template <typename T, int op> struct Opt {};\r
     \r
     template <typename T>\r
-    struct Cmp<T, MIN> \r
+    struct Opt<T, MIN> \r
     {\r
         static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval)\r
         {\r
@@ -422,7 +422,7 @@ namespace cv { namespace gpu { namespace mathfunc
     };\r
 \r
     template <typename T>\r
-    struct Cmp<T, MAX> \r
+    struct Opt<T, MAX> \r
     {\r
         static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval)\r
         {\r
@@ -448,23 +448,22 @@ namespace cv { namespace gpu { namespace mathfunc
 \r
         __syncthreads();\r
 \r
-        if (nthreads >= 512) if (tid < 256) { Cmp<best_type, op>::call(tid, 256, soptval); __syncthreads(); }\r
-        if (nthreads >= 256) if (tid < 128) { Cmp<best_type, op>::call(tid, 128, soptval); __syncthreads(); }\r
-        if (nthreads >= 128) if (tid < 64) { Cmp<best_type, op>::call(tid, 64, soptval); __syncthreads(); }\r
+        if (nthreads >= 512) if (tid < 256) { Opt<best_type, op>::call(tid, 256, soptval); __syncthreads(); }\r
+        if (nthreads >= 256) if (tid < 128) { Opt<best_type, op>::call(tid, 128, soptval); __syncthreads(); }\r
+        if (nthreads >= 128) if (tid < 64) { Opt<best_type, op>::call(tid, 64, soptval); __syncthreads(); }\r
 \r
         if (tid < 32)\r
         {\r
-            if (nthreads >= 64) Cmp<best_type, op>::call(tid, 32, soptval);\r
-            if (nthreads >= 32) Cmp<best_type, op>::call(tid, 16, soptval);\r
-            if (nthreads >= 16) Cmp<best_type, op>::call(tid, 8, soptval);\r
-            if (nthreads >= 8) Cmp<best_type, op>::call(tid, 4, soptval);\r
-            if (nthreads >= 4) Cmp<best_type, op>::call(tid, 2, soptval);\r
-            if (nthreads >= 2) Cmp<best_type, op>::call(tid, 1, soptval);\r
+            if (nthreads >= 64) Opt<best_type, op>::call(tid, 32, soptval);\r
+            if (nthreads >= 32) Opt<best_type, op>::call(tid, 16, soptval);\r
+            if (nthreads >= 16) Opt<best_type, op>::call(tid, 8, soptval);\r
+            if (nthreads >= 8) Opt<best_type, op>::call(tid, 4, soptval);\r
+            if (nthreads >= 4) Opt<best_type, op>::call(tid, 2, soptval);\r
+            if (nthreads >= 2) Opt<best_type, op>::call(tid, 1, soptval);\r
         }\r
 \r
         if (tid == 0) ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = (T)soptval[0];\r
     }\r
-\r
    \r
     template <typename T>\r
     void min_max_caller(const DevMem2D src, double* minval, double* maxval)\r
@@ -472,17 +471,19 @@ namespace cv { namespace gpu { namespace mathfunc
         dim3 threads(32, 8);\r
 \r
         // Allocate memory for aux. buffers\r
-        DevMem2D minval_buf[2]; DevMem2D maxval_buf[2];\r
+        DevMem2D minval_buf[2]; \r
         minval_buf[0].cols = divUp(src.cols, threads.x); \r
         minval_buf[0].rows = divUp(src.rows, threads.y);\r
         minval_buf[1].cols = divUp(minval_buf[0].cols, threads.x); \r
         minval_buf[1].rows = divUp(minval_buf[0].rows, threads.y);\r
+        cudaSafeCall(cudaMallocPitch(&minval_buf[0].data, &minval_buf[0].step, minval_buf[0].cols * sizeof(T), minval_buf[0].rows));\r
+        cudaSafeCall(cudaMallocPitch(&minval_buf[1].data, &minval_buf[1].step, minval_buf[1].cols * sizeof(T), minval_buf[1].rows));\r
+\r
+        DevMem2D maxval_buf[2];        \r
         maxval_buf[0].cols = divUp(src.cols, threads.x); \r
         maxval_buf[0].rows = divUp(src.rows, threads.y);\r
         maxval_buf[1].cols = divUp(maxval_buf[0].cols, threads.x); \r
         maxval_buf[1].rows = divUp(maxval_buf[0].rows, threads.y);\r
-        cudaSafeCall(cudaMallocPitch(&minval_buf[0].data, &minval_buf[0].step, minval_buf[0].cols * sizeof(T), minval_buf[0].rows));\r
-        cudaSafeCall(cudaMallocPitch(&minval_buf[1].data, &minval_buf[1].step, minval_buf[1].cols * sizeof(T), minval_buf[1].rows));\r
         cudaSafeCall(cudaMallocPitch(&maxval_buf[0].data, &maxval_buf[0].step, maxval_buf[0].cols * sizeof(T), maxval_buf[0].rows));\r
         cudaSafeCall(cudaMallocPitch(&maxval_buf[1].data, &maxval_buf[1].step, maxval_buf[1].cols * sizeof(T), maxval_buf[1].rows));\r
 \r
@@ -528,4 +529,219 @@ namespace cv { namespace gpu { namespace mathfunc
     template void min_max_caller<float>(const DevMem2D, double*, double*);\r
     template void min_max_caller<double>(const DevMem2D, double*, double*);\r
 \r
+    template <typename T, int op> struct OptLoc {};\r
+    \r
+    template <typename T>\r
+    struct OptLoc<T, MIN> \r
+    {\r
+        static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval, volatile unsigned int* optloc)\r
+        {\r
+            T val = optval[tid + offset];\r
+            if (val < optval[tid])\r
+            {\r
+                optval[tid] = val;\r
+                optloc[tid] = optloc[tid + offset];\r
+            }\r
+        }\r
+    };\r
+\r
+    template <typename T>\r
+    struct OptLoc<T, MAX> \r
+    {\r
+        static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval, volatile unsigned int* optloc)\r
+        {\r
+            T val = optval[tid + offset];\r
+            if (val > optval[tid])\r
+            {\r
+                optval[tid] = val;\r
+                optloc[tid] = optloc[tid + offset];\r
+            }\r
+        }\r
+    };\r
+\r
+    template <int nthreads, int op, typename T>\r
+    __global__ void opt_loc_init_kernel(int cols, int rows, const PtrStep src, PtrStep optval, PtrStep optloc)\r
+    {\r
+        typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
+        __shared__ best_type soptval[nthreads];\r
+        __shared__ unsigned int soptloc[nthreads];\r
+\r
+        unsigned int x0 = blockIdx.x * blockDim.x;\r
+        unsigned int y0 = blockIdx.y * blockDim.y;\r
+        unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
+\r
+        if (x0 + threadIdx.x < cols && y0 + threadIdx.y < rows)\r
+        {\r
+            soptval[tid] = ((const T*)src.ptr(y0 + threadIdx.y))[x0 + threadIdx.x];\r
+            soptloc[tid] = (y0 + threadIdx.y) * cols + x0 + threadIdx.x;\r
+        }\r
+        else\r
+        {\r
+            soptval[tid] = ((const T*)src.ptr(y0))[x0];\r
+            soptloc[tid] = y0 * cols + x0;\r
+        }\r
+\r
+        __syncthreads();\r
+\r
+        if (nthreads >= 512) if (tid < 256) { OptLoc<best_type, op>::call(tid, 256, soptval, soptloc); __syncthreads(); }\r
+        if (nthreads >= 256) if (tid < 128) { OptLoc<best_type, op>::call(tid, 128, soptval, soptloc); __syncthreads(); }\r
+        if (nthreads >= 128) if (tid < 64) { OptLoc<best_type, op>::call(tid, 64, soptval, soptloc); __syncthreads(); }\r
+\r
+        if (tid < 32)\r
+        {\r
+            if (nthreads >= 64) OptLoc<best_type, op>::call(tid, 32, soptval, soptloc);\r
+            if (nthreads >= 32) OptLoc<best_type, op>::call(tid, 16, soptval, soptloc);\r
+            if (nthreads >= 16) OptLoc<best_type, op>::call(tid, 8, soptval, soptloc);\r
+            if (nthreads >= 8) OptLoc<best_type, op>::call(tid, 4, soptval, soptloc);\r
+            if (nthreads >= 4) OptLoc<best_type, op>::call(tid, 2, soptval, soptloc);\r
+            if (nthreads >= 2) OptLoc<best_type, op>::call(tid, 1, soptval, soptloc);\r
+        }\r
+\r
+        if (tid == 0) \r
+        {\r
+            ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = (T)soptval[0];\r
+            ((unsigned int*)optloc.ptr(blockIdx.y))[blockIdx.x] = soptloc[0];\r
+        }\r
+    }\r
+\r
+    template <int nthreads, int op, typename T>\r
+    __global__ void opt_loc_kernel(int cols, int rows, const PtrStep src, const PtrStep loc, PtrStep optval, PtrStep optloc)\r
+    {\r
+        typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
+        __shared__ best_type soptval[nthreads];\r
+        __shared__ unsigned int soptloc[nthreads];\r
+\r
+        unsigned int x0 = blockIdx.x * blockDim.x;\r
+        unsigned int y0 = blockIdx.y * blockDim.y;\r
+        unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
+\r
+        if (x0 + threadIdx.x < cols && y0 + threadIdx.y < rows)\r
+        {\r
+            soptval[tid] = ((const T*)src.ptr(y0 + threadIdx.y))[x0 + threadIdx.x];\r
+            soptloc[tid] = ((const unsigned int*)loc.ptr(y0 + threadIdx.y))[x0 + threadIdx.x];\r
+        }\r
+        else\r
+        {\r
+            soptval[tid] = ((const T*)src.ptr(y0))[x0];\r
+            soptloc[tid] = ((const unsigned int*)loc.ptr(y0))[x0];\r
+        }\r
+\r
+        __syncthreads();\r
+\r
+        if (nthreads >= 512) if (tid < 256) { OptLoc<best_type, op>::call(tid, 256, soptval, soptloc); __syncthreads(); }\r
+        if (nthreads >= 256) if (tid < 128) { OptLoc<best_type, op>::call(tid, 128, soptval, soptloc); __syncthreads(); }\r
+        if (nthreads >= 128) if (tid < 64) { OptLoc<best_type, op>::call(tid, 64, soptval, soptloc); __syncthreads(); }\r
+\r
+        if (tid < 32)\r
+        {\r
+            if (nthreads >= 64) OptLoc<best_type, op>::call(tid, 32, soptval, soptloc);\r
+            if (nthreads >= 32) OptLoc<best_type, op>::call(tid, 16, soptval, soptloc);\r
+            if (nthreads >= 16) OptLoc<best_type, op>::call(tid, 8, soptval, soptloc);\r
+            if (nthreads >= 8) OptLoc<best_type, op>::call(tid, 4, soptval, soptloc);\r
+            if (nthreads >= 4) OptLoc<best_type, op>::call(tid, 2, soptval, soptloc);\r
+            if (nthreads >= 2) OptLoc<best_type, op>::call(tid, 1, soptval, soptloc);\r
+        }\r
+\r
+        if (tid == 0) \r
+        {\r
+            ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = (T)soptval[0];\r
+            ((unsigned int*)optloc.ptr(blockIdx.y))[blockIdx.x] = soptloc[0];\r
+        }\r
+    }\r
+\r
+    template <typename T>\r
+    void min_max_loc_caller(const DevMem2D src, double* minval, double* maxval, int* minlocx, int* minlocy, \r
+                                                                                int* maxlocx, int* maxlocy)\r
+    {\r
+        dim3 threads(32, 8);\r
+\r
+        // Allocate memory for aux. buffers\r
+\r
+        DevMem2D minval_buf[2]; \r
+        minval_buf[0].cols = divUp(src.cols, threads.x); \r
+        minval_buf[0].rows = divUp(src.rows, threads.y);\r
+        minval_buf[1].cols = divUp(minval_buf[0].cols, threads.x); \r
+        minval_buf[1].rows = divUp(minval_buf[0].rows, threads.y);\r
+        cudaSafeCall(cudaMallocPitch(&minval_buf[0].data, &minval_buf[0].step, minval_buf[0].cols * sizeof(T), minval_buf[0].rows));\r
+        cudaSafeCall(cudaMallocPitch(&minval_buf[1].data, &minval_buf[1].step, minval_buf[1].cols * sizeof(T), minval_buf[1].rows));\r
+\r
+        DevMem2D maxval_buf[2];        \r
+        maxval_buf[0].cols = divUp(src.cols, threads.x); \r
+        maxval_buf[0].rows = divUp(src.rows, threads.y);\r
+        maxval_buf[1].cols = divUp(maxval_buf[0].cols, threads.x); \r
+        maxval_buf[1].rows = divUp(maxval_buf[0].rows, threads.y);\r
+        cudaSafeCall(cudaMallocPitch(&maxval_buf[0].data, &maxval_buf[0].step, maxval_buf[0].cols * sizeof(T), maxval_buf[0].rows));\r
+        cudaSafeCall(cudaMallocPitch(&maxval_buf[1].data, &maxval_buf[1].step, maxval_buf[1].cols * sizeof(T), maxval_buf[1].rows));\r
+\r
+        DevMem2D minloc_buf[2]; \r
+        minloc_buf[0].cols = divUp(src.cols, threads.x); \r
+        minloc_buf[0].rows = divUp(src.rows, threads.y);\r
+        minloc_buf[1].cols = divUp(minloc_buf[0].cols, threads.x); \r
+        minloc_buf[1].rows = divUp(minloc_buf[0].rows, threads.y);\r
+        cudaSafeCall(cudaMallocPitch(&minloc_buf[0].data, &minloc_buf[0].step, minloc_buf[0].cols * sizeof(int), minloc_buf[0].rows));\r
+        cudaSafeCall(cudaMallocPitch(&minloc_buf[1].data, &minloc_buf[1].step, minloc_buf[1].cols * sizeof(int), minloc_buf[1].rows));\r
+\r
+        DevMem2D maxloc_buf[2]; \r
+        maxloc_buf[0].cols = divUp(src.cols, threads.x); \r
+        maxloc_buf[0].rows = divUp(src.rows, threads.y);\r
+        maxloc_buf[1].cols = divUp(maxloc_buf[0].cols, threads.x); \r
+        maxloc_buf[1].rows = divUp(maxloc_buf[0].rows, threads.y);\r
+        cudaSafeCall(cudaMallocPitch(&maxloc_buf[0].data, &maxloc_buf[0].step, maxloc_buf[0].cols * sizeof(int), maxloc_buf[0].rows));\r
+        cudaSafeCall(cudaMallocPitch(&maxloc_buf[1].data, &maxloc_buf[1].step, maxloc_buf[1].cols * sizeof(int), maxloc_buf[1].rows));\r
+\r
+        int curbuf = 0;\r
+        dim3 cursize(src.cols, src.rows);\r
+        dim3 grid(divUp(cursize.x, threads.x), divUp(cursize.y, threads.y));\r
+\r
+        opt_loc_init_kernel<256, MIN, T><<<grid, threads>>>(cursize.x, cursize.y, src, minval_buf[curbuf], minloc_buf[curbuf]);\r
+        opt_loc_init_kernel<256, MAX, T><<<grid, threads>>>(cursize.x, cursize.y, src, maxval_buf[curbuf], maxloc_buf[curbuf]);\r
+        cursize = grid;\r
+      \r
+        while (cursize.x > 1 || cursize.y > 1)\r
+        {\r
+            grid.x = divUp(cursize.x, threads.x); \r
+            grid.y = divUp(cursize.y, threads.y);  \r
+            opt_loc_kernel<256, MIN, T><<<grid, threads>>>(cursize.x, cursize.y, minval_buf[curbuf], minloc_buf[curbuf], \r
+                                                                                 minval_buf[1 - curbuf], minloc_buf[1 - curbuf]);\r
+            opt_loc_kernel<256, MAX, T><<<grid, threads>>>(cursize.x, cursize.y, maxval_buf[curbuf], maxloc_buf[curbuf], \r
+                                                                                 maxval_buf[1 - curbuf], maxloc_buf[1 - curbuf]);\r
+            curbuf = 1 - curbuf;\r
+            cursize = grid;\r
+        }\r
+\r
+        cudaSafeCall(cudaThreadSynchronize());\r
+\r
+        // Copy results from device to host\r
+\r
+        T minval_, maxval_;\r
+        cudaSafeCall(cudaMemcpy(&minval_, minval_buf[curbuf].ptr(0), sizeof(T), cudaMemcpyDeviceToHost));\r
+        cudaSafeCall(cudaMemcpy(&maxval_, maxval_buf[curbuf].ptr(0), sizeof(T), cudaMemcpyDeviceToHost));\r
+        *minval = minval_;\r
+        *maxval = maxval_;\r
+\r
+        unsigned int minloc, maxloc;\r
+        cudaSafeCall(cudaMemcpy(&minloc, minloc_buf[curbuf].ptr(0), sizeof(int), cudaMemcpyDeviceToHost));\r
+        cudaSafeCall(cudaMemcpy(&maxloc, maxloc_buf[curbuf].ptr(0), sizeof(int), cudaMemcpyDeviceToHost));\r
+        *minlocy = minloc / src.cols; *minlocx = minloc - *minlocy * src.cols;\r
+        *maxlocy = maxloc / src.cols; *maxlocx = maxloc - *maxlocy * src.cols;\r
+\r
+        // Release aux. buffers\r
+        cudaSafeCall(cudaFree(minval_buf[0].data));\r
+        cudaSafeCall(cudaFree(minval_buf[1].data));\r
+        cudaSafeCall(cudaFree(maxval_buf[0].data));\r
+        cudaSafeCall(cudaFree(maxval_buf[1].data));\r
+        cudaSafeCall(cudaFree(minloc_buf[0].data));\r
+        cudaSafeCall(cudaFree(minloc_buf[1].data));\r
+        cudaSafeCall(cudaFree(maxloc_buf[0].data));\r
+        cudaSafeCall(cudaFree(maxloc_buf[1].data));\r
+    }\r
+\r
+    template void min_max_loc_caller<unsigned char>(const DevMem2D, double*, double*, int*, int*, int*, int*);\r
+    template void min_max_loc_caller<signed char>(const DevMem2D, double*, double*, int*, int*, int*, int*);\r
+    template void min_max_loc_caller<unsigned short>(const DevMem2D, double*, double*, int*, int*, int*, int*);\r
+    template void min_max_loc_caller<signed short>(const DevMem2D, double*, double*, int*, int*, int*, int*);\r
+    template void min_max_loc_caller<int>(const DevMem2D, double*, double*, int*, int*, int*, int*);\r
+    template void min_max_loc_caller<float>(const DevMem2D, double*, double*, int*, int*, int*, int*);\r
+    template void min_max_loc_caller<double>(const DevMem2D, double*, double*, int*, int*, int*, int*);\r
+\r
 }}}\r

diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp
index f8b65ac..daabc9f 100644 (file)
--- a/tests/gpu/src/arithm.cpp
+++ b/tests/gpu/src/arithm.cpp
@@ -733,6 +733,71 @@ struct CV_GpuMinMaxTest: public CvTest
 };\r
 \r
 \r
+////////////////////////////////////////////////////////////////////////////////\r
+// Min max loc\r
+\r
+struct CV_GpuMinMaxLocTest: public CvTest\r
+{\r
+    CV_GpuMinMaxLocTest(): CvTest("GPU-MinMaxLocTest", "minMaxLoc") {}\r
+\r
+    void run(int)\r
+    {\r
+        for (int depth = CV_8U; depth <= CV_64F; ++depth)\r
+        {\r
+            int rows = 1, cols = 3;\r
+            test(rows, cols, depth);\r
+            for (int i = 0; i < 4; ++i)\r
+            {\r
+                int rows = 1 + rand() % 1000;\r
+                int cols = 1 + rand() % 1000;\r
+                test(rows, cols, depth);\r
+            }\r
+        }\r
+    }\r
+\r
+    void test(int rows, int cols, int depth)\r
+    {\r
+        cv::Mat src(rows, cols, depth);\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
+        double minVal, maxVal;\r
+        cv::Point minLoc, maxLoc;\r
+\r
+        if (depth != CV_8S)       \r
+            cv::minMaxLoc(src, &minVal, &maxVal, &minLoc, &maxLoc);\r
+        else \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
+            for (int i = 0; i < src.rows; ++i)\r
+                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
+                }\r
+        }\r
+\r
+        double minVal_, maxVal_;\r
+        cv::Point minLoc_, maxLoc_;        \r
+        cv::gpu::minMaxLoc(cv::gpu::GpuMat(src), &minVal_, &maxVal_, &minLoc_, &maxLoc_);\r
+       \r
+        CHECK(minVal == minVal_, CvTS::FAIL_INVALID_OUTPUT);\r
+        CHECK(maxVal == maxVal_, CvTS::FAIL_INVALID_OUTPUT);\r
+        CHECK(0 == memcmp(src.ptr(minLoc.y) + minLoc.x * src.elemSize(), src.ptr(minLoc_.y) + minLoc_.x * src.elemSize(), src.elemSize()),  \r
+              CvTS::FAIL_INVALID_OUTPUT);\r
+        CHECK(0 == memcmp(src.ptr(maxLoc.y) + maxLoc.x * src.elemSize(), src.ptr(maxLoc_.y) + maxLoc_.x * src.elemSize(), src.elemSize()),  \r
+              CvTS::FAIL_INVALID_OUTPUT);\r
+    }  \r
+};\r
+\r
+\r
 /////////////////////////////////////////////////////////////////////////////\r
 /////////////////// tests registration  /////////////////////////////////////\r
 /////////////////////////////////////////////////////////////////////////////\r
@@ -760,3 +825,4 @@ CV_GpuNppImagePhaseTest CV_GpuNppImagePhase_test;
 CV_GpuNppImageCartToPolarTest CV_GpuNppImageCartToPolar_test;\r
 CV_GpuNppImagePolarToCartTest CV_GpuNppImagePolarToCart_test;\r
 CV_GpuMinMaxTest CV_GpuMinMaxTest_test;\r
+CV_GpuMinMaxLocTest CV_GpuMinMaxLocTest_test;\r