This source file includes following definitions.
- initialValue
- result
- initialValue
- result
- initialValue
- result
- initialValue
- result
- gridReduceToRow
- gridReduceToRow
- gridReduceToColumn_
- gridReduceToColumn_
- gridReduceToColumn
- gridReduceToColumn
#pragma once
#ifndef __OPENCV_CUDEV_GRID_REDUCE_TO_VEC_HPP__
#define __OPENCV_CUDEV_GRID_REDUCE_TO_VEC_HPP__
#include "../common.hpp"
#include "../util/vec_traits.hpp"
#include "../util/limits.hpp"
#include "../util/saturate_cast.hpp"
#include "../ptr2d/traits.hpp"
#include "../ptr2d/gpumat.hpp"
#include "../ptr2d/mask.hpp"
#include "../functional/functional.hpp"
#include "detail/reduce_to_column.hpp"
#include "detail/reduce_to_row.hpp"
namespace cv { namespace cudev {
template <typename T> struct Sum : plus<T>
{
typedef T work_type;
template <typename U> struct rebind
{
typedef Sum<U> other;
};
__device__ __forceinline__ static T initialValue()
{
return VecTraits<T>::all(0);
}
__device__ __forceinline__ static T result(T r, int)
{
return r;
}
};
template <typename T> struct Avg : plus<T>
{
typedef T work_type;
template <typename U> struct rebind
{
typedef Avg<U> other;
};
__device__ __forceinline__ static T initialValue()
{
return VecTraits<T>::all(0);
}
__device__ __forceinline__ static T result(T r, float sz)
{
return saturate_cast<T>(r / sz);
}
};
template <typename T> struct Min : minimum<T>
{
typedef T work_type;
template <typename U> struct rebind
{
typedef Min<U> other;
};
__device__ __forceinline__ static T initialValue()
{
return VecTraits<T>::all(numeric_limits<typename VecTraits<T>::elem_type>::max());
}
__device__ __forceinline__ static T result(T r, int)
{
return r;
}
};
template <typename T> struct Max : maximum<T>
{
typedef T work_type;
template <typename U> struct rebind
{
typedef Max<U> other;
};
__device__ __forceinline__ static T initialValue()
{
return VecTraits<T>::all(-numeric_limits<typename VecTraits<T>::elem_type>::max());
}
__device__ __forceinline__ static T result(T r, int)
{
return r;
}
};
template <class Reductor, class SrcPtr, typename ResType, class MaskPtr>
__host__ void gridReduceToRow(const SrcPtr& src, GpuMat_<ResType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
dst.create(1, cols);
grid_reduce_to_vec_detail::reduceToRow<Reductor>(shrinkPtr(src),
dst[0],
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Reductor, class SrcPtr, typename ResType>
__host__ void gridReduceToRow(const SrcPtr& src, GpuMat_<ResType>& dst, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
dst.create(1, cols);
grid_reduce_to_vec_detail::reduceToRow<Reductor>(shrinkPtr(src),
dst[0],
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Reductor, class Policy, class SrcPtr, typename ResType, class MaskPtr>
__host__ void gridReduceToColumn_(const SrcPtr& src, GpuMat_<ResType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
dst.create(1, rows);
grid_reduce_to_vec_detail::reduceToColumn<Reductor, Policy>(shrinkPtr(src),
dst[0],
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Reductor, class Policy, class SrcPtr, typename ResType>
__host__ void gridReduceToColumn_(const SrcPtr& src, GpuMat_<ResType>& dst, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
dst.create(1, rows);
grid_reduce_to_vec_detail::reduceToColumn<Reductor, Policy>(shrinkPtr(src),
dst[0],
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
struct DefaultReduceToVecPolicy
{
enum {
block_size_x = 32,
block_size_y = 8
};
};
template <class Reductor, class SrcPtr, typename ResType, class MaskPtr>
__host__ void gridReduceToColumn(const SrcPtr& src, GpuMat_<ResType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridReduceToColumn_<Reductor, DefaultReduceToVecPolicy>(src, dst, mask, stream);
}
template <class Reductor, class SrcPtr, typename ResType>
__host__ void gridReduceToColumn(const SrcPtr& src, GpuMat_<ResType>& dst, Stream& stream = Stream::Null())
{
gridReduceToColumn_<Reductor, DefaultReduceToVecPolicy>(src, dst, stream);
}
}}
#endif