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
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
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
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
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
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
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
}\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
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
\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
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
\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
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