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::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::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
{\r
using namespace mathfunc::minmax;\r
\r
- double maxVal_;\r
- if (!maxVal) maxVal = &maxVal_;\r
+ double minVal_; if (!minVal) minVal = &minVal_;\r
+ double maxVal_; if (!maxVal) maxVal = &maxVal_;\r
\r
GpuMat src_ = src.reshape(1);\r
\r
\r
namespace cv { namespace gpu { namespace mathfunc { namespace minmaxloc {\r
\r
+ void get_buf_size_required(int elem_size, int& b1cols, int& b1rows, \r
+ int& b2cols, int& b2rows);\r
+\r
template <typename T> \r
void min_max_loc_caller(const DevMem2D src, double* minval, double* maxval, \r
- int* minlocx, int* minlocy, int* maxlocx, int* maxlocy);\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
+ int minloc[2], int maxloc[2], PtrStep valbuf, PtrStep locbuf);\r
\r
}}}}\r
\r
void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc)\r
+{ \r
+ GpuMat valbuf, locbuf;\r
+ minMaxLoc(src, minVal, maxVal, minLoc, maxLoc, 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
{\r
using namespace mathfunc::minmaxloc;\r
-\r
CV_Assert(src.channels() == 1);\r
\r
- double maxVal_;\r
- if (!maxVal) maxVal = &maxVal_;\r
+ double minVal_; if (!minVal) minVal = &minVal_;\r
+ double maxVal_; if (!maxVal) maxVal = &maxVal_;\r
+ int minLoc_[2];\r
+ int maxLoc_[2];\r
\r
- cv::Point minLoc_;\r
- if (!minLoc) minLoc = &minLoc_;\r
+ Size valbuf_size, locbuf_size;\r
+ get_buf_size_required(src.elemSize(), valbuf_size.width, valbuf_size.height, \r
+ locbuf_size.width, locbuf_size.height);\r
+ valbuf.create(valbuf_size, CV_8U);\r
+ locbuf.create(locbuf_size, CV_8U);\r
\r
- cv::Point maxLoc_;\r
- if (!maxLoc) maxLoc = &maxLoc_;\r
- \r
- switch (src.type())\r
+ int major, minor;\r
+ getComputeCapability(getDevice(), major, minor);\r
+ \r
+ if (major >= 1 && minor >= 1)\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_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_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: min_max_loc_caller<double>(src, minVal, maxVal, minLoc_, maxLoc_, valbuf, locbuf); break;\r
+ default: CV_Error(CV_StsBadArg, "Unsupported type");\r
+ }\r
+ }\r
+ else\r
{\r
- case CV_8U:\r
- min_max_loc_caller<unsigned char>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
- break;\r
- case CV_8S:\r
- min_max_loc_caller<signed char>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
- break;\r
- case CV_16U:\r
- min_max_loc_caller<unsigned short>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
- break;\r
- case CV_16S:\r
- min_max_loc_caller<signed short>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
- break;\r
- case CV_32S:\r
- min_max_loc_caller<int>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
- break;\r
- case CV_32F:\r
- min_max_loc_caller<float>(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y);\r
- break;\r
- case CV_64F:\r
- 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
+ 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
+ default: CV_Error(CV_StsBadArg, "Unsupported type");\r
+ }\r
}\r
+\r
+ if (minLoc) { minLoc->x = minLoc_[0]; minLoc->y = minLoc_[1]; }\r
+ if (maxLoc) { maxLoc->x = maxLoc_[0]; maxLoc->y = maxLoc_[1]; }\r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
template <> struct MinMaxTypeTraits<float> { typedef float best_type; };\r
template <> struct MinMaxTypeTraits<double> { typedef double best_type; };\r
\r
- // Available optimization operations\r
- enum { OP_MIN, OP_MAX };\r
\r
namespace minmax \r
{\r
\r
\r
template <int nthreads, typename T>\r
- __global__ void min_max_kernel(int cols, int rows, const PtrStep src, T* minval, T* maxval)\r
+ __global__ void min_max_kernel(const DevMem2D src, T* minval, T* maxval)\r
{\r
typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
__shared__ best_type sminval[nthreads];\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 < rows; ++y)\r
+ for (unsigned int y = 0; y < ctheight && y0 + y * blockDim.y < src.rows; ++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 < cols; ++x)\r
+ for (unsigned int x = 0; x < ctwidth && x0 + x * blockDim.x < src.cols; ++x)\r
{\r
val = ptr[x0 + x * blockDim.x];\r
mymin = min(mymin, val);\r
if (nthreads >= 2) merge(tid, 1, sminval, smaxval);\r
}\r
\r
- __syncthreads();\r
-\r
if (tid == 0) \r
{\r
minval[blockIdx.y * gridDim.x + blockIdx.x] = (T)sminval[0];\r
__threadfence();\r
if (atomicInc(&blocks_finished, gridDim.x * gridDim.y) == gridDim.x * gridDim.y - 1)\r
{\r
- mymin = numeric_limits_gpu<T>::max();\r
- mymax = numeric_limits_gpu<T>::min();\r
- for (unsigned int i = 0; i < gridDim.x * gridDim.y; ++i)\r
+ mymin = minval[0];\r
+ mymax = maxval[0];\r
+ for (unsigned int i = 1; i < gridDim.x * gridDim.y; ++i)\r
{ \r
mymin = min(mymin, minval[i]);\r
mymax = max(mymax, maxval[i]);\r
T* maxval_buf = (T*)buf.ptr(1);\r
\r
cudaSafeCall(cudaMemcpyToSymbol(blocks_finished, &czero, sizeof(blocks_finished)));\r
- min_max_kernel<256, T><<<grid, threads>>>(src.cols, src.rows, src, minval_buf, maxval_buf);\r
+ min_max_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf);\r
cudaSafeCall(cudaThreadSynchronize());\r
\r
T minval_, maxval_;\r
__global__ void min_max_kernel_2ndstep(T* minval, T* maxval, int size)\r
{\r
T val;\r
- T mymin = numeric_limits_gpu<T>::max();\r
- T mymax = numeric_limits_gpu<T>::min();\r
- for (unsigned int i = 0; i < size; ++i)\r
+ T mymin = minval[0];\r
+ T mymax = maxval[0];\r
+ for (unsigned int i = 1; i < size; ++i)\r
{ \r
val = minval[i]; if (val < mymin) mymin = val;\r
val = maxval[i]; if (val > mymax) mymax = val;\r
T* maxval_buf = (T*)buf.ptr(1);\r
\r
cudaSafeCall(cudaMemcpyToSymbol(blocks_finished, &czero, sizeof(blocks_finished)));\r
- min_max_kernel<256, T><<<grid, threads>>>(src.cols, src.rows, src, minval_buf, maxval_buf);\r
+ min_max_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf);\r
min_max_kernel_2ndstep<T><<<1, 1>>>(minval_buf, maxval_buf, grid.x * grid.y);\r
cudaSafeCall(cudaThreadSynchronize());\r
\r
\r
namespace minmaxloc {\r
\r
- template <typename T, int op> struct OptLoc {};\r
- \r
- template <typename T>\r
- struct OptLoc<T, OP_MIN> \r
+ __constant__ int ctwidth;\r
+ __constant__ int ctheight;\r
+\r
+ static const unsigned int czero = 0;\r
+\r
+ // Global counter of blocks finished its work\r
+ __device__ unsigned int blocks_finished;\r
+\r
+\r
+ // Estimates good thread configuration\r
+ // - threads variable satisfies to threads.x * threads.y == 256\r
+ void estimate_thread_cfg(dim3& threads, dim3& grid)\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
+ threads = dim3(64, 4);\r
+ grid = dim3(6, 5);\r
+ }\r
+\r
+\r
+ // Returns required buffer sizes\r
+ void get_buf_size_required(int elem_size, int& b1cols, int& b1rows, \r
+ int& b2cols, int& b2rows)\r
+ {\r
+ dim3 threads, grid;\r
+ estimate_thread_cfg(threads, grid);\r
+ b1cols = grid.x * grid.y * elem_size; // For values\r
+ b1rows = 2;\r
+ b2cols = grid.x * grid.y * sizeof(int); // For locations\r
+ b2rows = 2;\r
+ }\r
+\r
+\r
+ // Estimates device constants which are used in the kernels using specified thread configuration\r
+ void estimate_kernel_consts(int cols, int rows, const dim3& threads, const dim3& grid)\r
+ { \r
+ int twidth = divUp(divUp(cols, grid.x), threads.x);\r
+ int theight = divUp(divUp(rows, grid.y), threads.y);\r
+ cudaSafeCall(cudaMemcpyToSymbol(ctwidth, &twidth, sizeof(ctwidth))); \r
+ cudaSafeCall(cudaMemcpyToSymbol(ctheight, &theight, sizeof(ctheight))); \r
+ } \r
+\r
\r
template <typename T>\r
- struct OptLoc<T, OP_MAX> \r
+ __device__ void merge(unsigned int tid, unsigned int offset, volatile T* minval, volatile T* maxval, \r
+ volatile unsigned int* minloc, volatile unsigned int* maxloc)\r
{\r
- static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval, volatile unsigned int* optloc)\r
+ T val = minval[tid + offset];\r
+ if (val < minval[tid])\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
+ minval[tid] = val;\r
+ minloc[tid] = minloc[tid + offset];\r
}\r
- };\r
+ val = maxval[tid + offset];\r
+ if (val > maxval[tid])\r
+ {\r
+ maxval[tid] = val;\r
+ maxloc[tid] = maxloc[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
+ template <int nthreads, typename T>\r
+ __global__ void min_max_loc_kernel(const DevMem2D src, T* minval, T* maxval, \r
+ unsigned int* minloc, unsigned int* maxloc)\r
{\r
typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
- __shared__ best_type soptval[nthreads];\r
- __shared__ unsigned int soptloc[nthreads];\r
+ __shared__ best_type sminval[nthreads];\r
+ __shared__ best_type smaxval[nthreads];\r
+ __shared__ unsigned int sminloc[nthreads];\r
+ __shared__ unsigned int smaxloc[nthreads];\r
\r
- unsigned int x0 = blockIdx.x * blockDim.x;\r
- unsigned int y0 = blockIdx.y * blockDim.y;\r
+ unsigned int x0 = blockIdx.x * blockDim.x * ctwidth + threadIdx.x;\r
+ unsigned int y0 = blockIdx.y * blockDim.y * ctheight + threadIdx.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
+ T val = ((const T*)src.ptr(0))[0];\r
+ T mymin = val, mymax = val; \r
+ unsigned int myminloc = 0, mymaxloc = 0;\r
+ for (unsigned int y = 0; y < ctheight && y0 + y * blockDim.y < src.rows; ++y)\r
{\r
- soptval[tid] = ((const T*)src.ptr(y0))[x0];\r
- soptloc[tid] = y0 * cols + x0;\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
+ {\r
+ val = ptr[x0 + x * blockDim.x];\r
+ if (val < mymin) \r
+ { \r
+ mymin = val; \r
+ myminloc = (y0 + y * blockDim.y) * src.cols + x0 + x * blockDim.x; \r
+ }\r
+ else if (val > mymax)\r
+ {\r
+ mymax = val; \r
+ mymaxloc = (y0 + y * blockDim.y) * src.cols + x0 + x * blockDim.x; \r
+ }\r
+ }\r
}\r
\r
+ sminval[tid] = mymin; \r
+ smaxval[tid] = mymax;\r
+ sminloc[tid] = myminloc;\r
+ smaxloc[tid] = mymaxloc;\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
+ if (nthreads >= 512) if (tid < 256) { merge(tid, 256, sminval, smaxval, sminloc, smaxloc); __syncthreads(); }\r
+ if (nthreads >= 256) if (tid < 128) { merge(tid, 128, sminval, smaxval, sminloc, smaxloc); __syncthreads(); }\r
+ if (nthreads >= 128) if (tid < 64) { merge(tid, 64, sminval, smaxval, sminloc, smaxloc); __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
+ if (nthreads >= 64) merge(tid, 32, sminval, smaxval, sminloc, smaxloc);\r
+ if (nthreads >= 32) merge(tid, 16, sminval, smaxval, sminloc, smaxloc);\r
+ if (nthreads >= 16) merge(tid, 8, sminval, smaxval, sminloc, smaxloc);\r
+ if (nthreads >= 8) merge(tid, 4, sminval, smaxval, sminloc, smaxloc);\r
+ if (nthreads >= 4) merge(tid, 2, sminval, smaxval, sminloc, smaxloc);\r
+ if (nthreads >= 2) merge(tid, 1, sminval, smaxval, sminloc, smaxloc);\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
+ minval[blockIdx.y * gridDim.x + blockIdx.x] = (T)sminval[0];\r
+ maxval[blockIdx.y * gridDim.x + blockIdx.x] = (T)smaxval[0];\r
+ minloc[blockIdx.y * gridDim.x + blockIdx.x] = sminloc[0];\r
+ maxloc[blockIdx.y * gridDim.x + blockIdx.x] = smaxloc[0];\r
+ }\r
+\r
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 110\r
+ \r
+ // Process partial results in the first thread of the last block \r
+ if ((gridDim.x > 1 || gridDim.y > 1) && tid == 0)\r
+ {\r
+ __threadfence();\r
+ if (atomicInc(&blocks_finished, gridDim.x * gridDim.y) == gridDim.x * gridDim.y - 1)\r
+ {\r
+ mymin = minval[0];\r
+ mymax = maxval[0];\r
+ unsigned int imin = 0, imax = 0;\r
+ for (unsigned int i = 1; i < gridDim.x * gridDim.y; ++i)\r
+ { \r
+ val = minval[i]; if (val < mymin) { mymin = val; imin = i; }\r
+ val = maxval[i]; if (val > mymax) { mymax = val; imax = i; }\r
+ }\r
+ minval[0] = mymin;\r
+ maxval[0] = mymax;\r
+ minloc[0] = minloc[imin];\r
+ maxloc[0] = maxloc[imax];\r
+ }\r
}\r
+#endif\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
+ 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
- typedef typename MinMaxTypeTraits<T>::best_type best_type;\r
- __shared__ best_type soptval[nthreads];\r
- __shared__ unsigned int soptloc[nthreads];\r
+ dim3 threads, grid;\r
+ estimate_thread_cfg(threads, grid);\r
+ estimate_kernel_consts(src.cols, src.rows, threads, grid);\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
+ 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
- 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
+ cudaSafeCall(cudaMemcpyToSymbol(blocks_finished, &czero, sizeof(blocks_finished)));\r
+ min_max_loc_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf, minloc_buf, maxloc_buf);\r
+ cudaSafeCall(cudaThreadSynchronize());\r
\r
- __syncthreads();\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
- 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
+ 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
- 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
+ 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<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<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
- 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
+ // This kernel will be used only when compute capability is 1.0\r
+ template <typename T>\r
+ __global__ void min_max_loc_kernel_2ndstep(T* minval, T* maxval, unsigned int* minloc, unsigned int* maxloc, int size)\r
+ {\r
+ T val;\r
+ T mymin = minval[0];\r
+ T mymax = maxval[0];\r
+ unsigned int imin = 0, imax = 0;\r
+ for (unsigned int i = 1; i < size; ++i)\r
+ { \r
+ val = minval[i]; if (val < mymin) { mymin = val; imin = i; }\r
+ val = maxval[i]; if (val > mymax) { mymax = val; imax = i; }\r
}\r
+ minval[0] = mymin;\r
+ maxval[0] = mymax;\r
+ minloc[0] = minloc[imin];\r
+ maxloc[0] = maxloc[imax];\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, OP_MIN, T><<<grid, threads>>>(cursize.x, cursize.y, src, minval_buf[curbuf], minloc_buf[curbuf]);\r
- opt_loc_init_kernel<256, OP_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, OP_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, OP_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
+ void min_max_loc_caller_2steps(const DevMem2D src, 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
- cudaSafeCall(cudaThreadSynchronize());\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
- // Copy results from device to host\r
+ cudaSafeCall(cudaMemcpyToSymbol(blocks_finished, &czero, sizeof(blocks_finished)));\r
+ min_max_loc_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf, minloc_buf, maxloc_buf);\r
+ min_max_loc_kernel_2ndstep<T><<<1, 1>>>(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[curbuf].ptr(0), sizeof(T), cudaMemcpyDeviceToHost));\r
- cudaSafeCall(cudaMemcpy(&maxval_, maxval_buf[curbuf].ptr(0), sizeof(T), cudaMemcpyDeviceToHost));\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[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
+ 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_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
\r
} // namespace minmaxloc\r
\r
projects / platform / upstream / opencv.git / commitdiff