#include "opencv2/gpu/device/vec_math.hpp"\r
#include "opencv2/gpu/device/saturate_cast.hpp"\r
#include "opencv2/gpu/device/filters.hpp"\r
+# include <cfloat>\r
\r
namespace cv { namespace gpu { namespace device\r
{\r
}\r
}\r
\r
+ template <typename Ptr2D, typename T> __global__ void resize_area(const Ptr2D src, float fx, float fy, DevMem2D_<T> dst)\r
+ {\r
+ const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+ const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+ if (x < dst.cols && y < dst.rows)\r
+ {\r
+ dst(y, x) = saturate_cast<T>(src(y, x));\r
+ }\r
+ }\r
+\r
template <template <typename> class Filter, typename T> struct ResizeDispatcherStream\r
{\r
static void call(DevMem2D_<T> src, float fx, float fy, DevMem2D_<T> dst, cudaStream_t stream)\r
\r
BrdReplicate<T> brd(src.rows, src.cols);\r
BorderReader< PtrStep<T>, BrdReplicate<T> > brdSrc(src, brd);\r
- Filter< BorderReader< PtrStep<T>, BrdReplicate<T> > > filteredSrc(brdSrc);\r
+ Filter< BorderReader< PtrStep<T>, BrdReplicate<T> > > filteredSrc(brdSrc, fx, fy);\r
\r
resize<<<grid, block, 0, stream>>>(filteredSrc, fx, fy, dst);\r
cudaSafeCall( cudaGetLastError() );\r
}\r
};\r
\r
+ template <typename T> struct ResizeDispatcherStream<AreaFilter, T>\r
+ {\r
+ static void call(DevMem2D_<T> src, float fx, float fy, DevMem2D_<T> dst, cudaStream_t stream)\r
+ {\r
+ dim3 block(32, 8);\r
+ dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));\r
+\r
+ BrdConstant<T> brd(src.rows, src.cols);\r
+ BorderReader< PtrStep<T>, BrdConstant<T> > brdSrc(src, brd);\r
+ AreaFilter< BorderReader< PtrStep<T>, BrdConstant<T> > > filteredSrc(brdSrc, fx, fy);\r
+ resize_area<<<grid, block, 0, stream>>>(filteredSrc, fx, fy, dst);\r
+ cudaSafeCall( cudaGetLastError() );\r
+ }\r
+ };\r
+\r
+ template <typename T> struct ResizeDispatcherStream<IntegerAreaFilter, T>\r
+ {\r
+ static void call(DevMem2D_<T> src, float fx, float fy, DevMem2D_<T> dst, cudaStream_t stream)\r
+ {\r
+ dim3 block(32, 8);\r
+ dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));\r
+\r
+ BrdConstant<T> brd(src.rows, src.cols);\r
+ BorderReader< PtrStep<T>, BrdConstant<T> > brdSrc(src, brd);\r
+ IntegerAreaFilter< BorderReader< PtrStep<T>, BrdConstant<T> > > filteredSrc(brdSrc, fx, fy);\r
+ resize_area<<<grid, block, 0, stream>>>(filteredSrc, fx, fy, dst);\r
+ cudaSafeCall( cudaGetLastError() );\r
+ }\r
+ };\r
+\r
template <template <typename> class Filter, typename T> struct ResizeDispatcherNonStream\r
{\r
static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, float fx, float fy, DevMem2D_<T> dst)\r
}\r
};\r
\r
+ template <typename T> struct ResizeDispatcher<AreaFilter, T>\r
+ {\r
+ static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, float fx, float fy, DevMem2D_<T> dst, cudaStream_t stream)\r
+ {\r
+ int iscale_x = round(fx);\r
+ int iscale_y = round(fy);\r
+\r
+ if( std::abs(fx - iscale_x) < FLT_MIN && std::abs(fy - iscale_y) < FLT_MIN)\r
+ ResizeDispatcherStream<IntegerAreaFilter, T>::call(src, fx, fy, dst, stream);\r
+ else\r
+ ResizeDispatcherStream<AreaFilter, T>::call(src, fx, fy, dst, stream);\r
+ }\r
+ };\r
+\r
template <typename T> void resize_gpu(DevMem2Db src, DevMem2Db srcWhole, int xoff, int yoff, float fx, float fy, \r
DevMem2Db dst, int interpolation, cudaStream_t stream)\r
{\r
typedef void (*caller_t)(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, float fx, float fy, DevMem2D_<T> dst, cudaStream_t stream);\r
\r
- static const caller_t callers[3] = \r
+ static const caller_t callers[4] =\r
{\r
- ResizeDispatcher<PointFilter, T>::call, ResizeDispatcher<LinearFilter, T>::call, ResizeDispatcher<CubicFilter, T>::call\r
+ ResizeDispatcher<PointFilter, T>::call,\r
+ ResizeDispatcher<LinearFilter, T>::call,\r
+ ResizeDispatcher<CubicFilter, T>::call,\r
+ ResizeDispatcher<AreaFilter, T>::call\r
};\r
\r
callers[interpolation](static_cast< DevMem2D_<T> >(src), static_cast< DevMem2D_<T> >(srcWhole), xoff, yoff, fx, fy, \r
typedef typename Ptr2D::elem_type elem_type;\r
typedef float index_type;\r
\r
- explicit __host__ __device__ __forceinline__ PointFilter(const Ptr2D& src_) : src(src_) {}\r
+ explicit __host__ __device__ __forceinline__ PointFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f) : src(src_) {}\r
\r
__device__ __forceinline__ elem_type operator ()(float y, float x) const\r
{\r
typedef typename Ptr2D::elem_type elem_type;\r
typedef float index_type;\r
\r
- explicit __host__ __device__ __forceinline__ LinearFilter(const Ptr2D& src_) : src(src_) {}\r
+ explicit __host__ __device__ __forceinline__ LinearFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f) : src(src_) {}\r
\r
__device__ __forceinline__ elem_type operator ()(float y, float x) const\r
{\r
typedef float index_type;\r
typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;\r
\r
- explicit __host__ __device__ __forceinline__ CubicFilter(const Ptr2D& src_) : src(src_) {}\r
+ explicit __host__ __device__ __forceinline__ CubicFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f) : src(src_) {}\r
\r
static __device__ __forceinline__ float bicubicCoeff(float x_)\r
{\r
\r
const Ptr2D src;\r
};\r
+ // for integer scaling\r
+ template <typename Ptr2D> struct IntegerAreaFilter\r
+ {\r
+ typedef typename Ptr2D::elem_type elem_type;\r
+ typedef float index_type;\r
+\r
+ explicit __host__ __device__ __forceinline__ IntegerAreaFilter(const Ptr2D& src_, float scale_x_, float scale_y_)\r
+ : src(src_), scale_x(scale_x_), scale_y(scale_y_), scale(1.f / (scale_x * scale_y)) {}\r
+\r
+ __device__ __forceinline__ elem_type operator ()(float y, float x) const\r
+ {\r
+ float fsx1 = x * scale_x;\r
+ float fsx2 = fsx1 + scale_x;\r
+\r
+ int sx1 = __float2int_ru(fsx1);\r
+ int sx2 = __float2int_rd(fsx2);\r
+\r
+ float fsy1 = y * scale_y;\r
+ float fsy2 = fsy1 + scale_y;\r
+\r
+ int sy1 = __float2int_ru(fsy1);\r
+ int sy2 = __float2int_rd(fsy2);\r
+\r
+ typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;\r
+ work_type out = VecTraits<work_type>::all(0.f);\r
+\r
+ for(int dy = sy1; dy < sy2; ++dy)\r
+ for(int dx = sx1; dx < sx2; ++dx)\r
+ {\r
+ out = out + src(dy, dx) * scale;\r
+ }\r
+\r
+ return saturate_cast<elem_type>(out);\r
+ }\r
+\r
+ const Ptr2D src;\r
+ float scale_x, scale_y ,scale;\r
+ };\r
+\r
+ template <typename Ptr2D> struct AreaFilter\r
+ {\r
+ typedef typename Ptr2D::elem_type elem_type;\r
+ typedef float index_type;\r
+\r
+ explicit __host__ __device__ __forceinline__ AreaFilter(const Ptr2D& src_, float scale_x_, float scale_y_)\r
+ : src(src_), scale_x(scale_x_), scale_y(scale_y_){}\r
+\r
+ __device__ __forceinline__ elem_type operator ()(float y, float x) const\r
+ {\r
+ float fsx1 = x * scale_x;\r
+ float fsx2 = fsx1 + scale_x;\r
+\r
+ int sx1 = __float2int_ru(fsx1);\r
+ int sx2 = __float2int_rd(fsx2);\r
+\r
+ float fsy1 = y * scale_y;\r
+ float fsy2 = fsy1 + scale_y;\r
+\r
+ int sy1 = __float2int_ru(fsy1);\r
+ int sy2 = __float2int_rd(fsy2);\r
+\r
+ float scale = 1.f / (fminf(scale_x, src.width - fsx1) * fminf(scale_y, src.height - fsy1));\r
+\r
+ typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;\r
+ work_type out = VecTraits<work_type>::all(0.f);\r
+\r
+ for (int dy = sy1; dy < sy2; ++dy)\r
+ {\r
+ for (int dx = sx1; dx < sx2; ++dx)\r
+ out = out + src(dy, dx) * scale;\r
+\r
+ if (sx1 > fsx1)\r
+ out = out + src(dy, (sx1 -1) ) * ((sx1 - fsx1) * scale);\r
+\r
+ if (sx2 < fsx2)\r
+ out = out + src(dy, sx2) * ((fsx2 -sx2) * scale);\r
+ }\r
+\r
+ if (sy1 > fsy1)\r
+ for (int dx = sx1; dx < sx2; ++dx)\r
+ out = out + src( (sy1 - 1) , dx) * ((sy1 -fsy1) * scale);\r
+\r
+ if (sy2 < fsy2)\r
+ for (int dx = sx1; dx < sx2; ++dx)\r
+ out = out + src(sy2, dx) * ((fsy2 -sy2) * scale);\r
+\r
+ if ((sy1 > fsy1) && (sx1 > fsx1))\r
+ out = out + src( (sy1 - 1) , (sx1 - 1)) * ((sy1 -fsy1) * (sx1 -fsx1) * scale);\r
+\r
+ if ((sy1 > fsy1) && (sx2 < fsx2))\r
+ out = out + src( (sy1 - 1) , sx2) * ((sy1 -fsy1) * (fsx2 -sx2) * scale);\r
+\r
+ if ((sy2 < fsy2) && (sx2 < fsx2))\r
+ out = out + src(sy2, sx2) * ((fsy2 -sy2) * (fsx2 -sx2) * scale);\r
+\r
+ if ((sy2 < fsy2) && (sx1 > fsx1))\r
+ out = out + src(sy2, (sx1 - 1)) * ((fsy2 -sy2) * (sx1 -fsx1) * scale);\r
+\r
+ return saturate_cast<elem_type>(out);\r
+ }\r
+\r
+ const Ptr2D src;\r
+ float scale_x, scale_y;\r
+ int width, haight;\r
+ };\r
}}} // namespace cv { namespace gpu { namespace device\r
\r
#endif // __OPENCV_GPU_FILTERS_HPP__\r
\r
template<> struct VecTraits<char> \r
{ \r
- typedef char elem_type; \r
+ typedef char elem_type;\r
enum {cn=1}; \r
static __device__ __host__ __forceinline__ char all(char v) {return v;}\r
static __device__ __host__ __forceinline__ char make(char x) {return x;}\r
};\r
template<> struct VecTraits<schar> \r
{ \r
- typedef schar elem_type; \r
+ typedef schar elem_type;\r
enum {cn=1}; \r
static __device__ __host__ __forceinline__ schar all(schar v) {return v;}\r
static __device__ __host__ __forceinline__ schar make(schar x) {return x;}\r
void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation, Stream& s)\r
{\r
CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);\r
- CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
+ CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR\r
+ || interpolation == INTER_CUBIC || interpolation == INTER_AREA);\r
CV_Assert(!(dsize == Size()) || (fx > 0 && fy > 0));\r
\r
if (dsize == Size())\r
\r
namespace\r
{\r
- template <typename T, template <typename> class Interpolator> void resizeImpl(const cv::Mat& src, cv::Mat& dst, double fx, double fy)\r
+ template <typename T, template <typename> class Interpolator>\r
+ void resizeImpl(const cv::Mat& src, cv::Mat& dst, double fx, double fy)\r
{\r
const int cn = src.channels();\r
\r
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)),\r
WHOLE_SUBMAT));\r
\r
+\r
+/////////////////\r
+PARAM_TEST_CASE(ResizeArea, cv::gpu::DeviceInfo, cv::Size, MatType, double, Interpolation, UseRoi)\r
+{\r
+ cv::gpu::DeviceInfo devInfo;\r
+ cv::Size size;\r
+ double coeff;\r
+ int interpolation;\r
+ int type;\r
+ bool useRoi;\r
+\r
+ virtual void SetUp()\r
+ {\r
+ devInfo = GET_PARAM(0);\r
+ size = GET_PARAM(1);\r
+ type = GET_PARAM(2);\r
+ coeff = GET_PARAM(3);\r
+ interpolation = GET_PARAM(4);\r
+ useRoi = GET_PARAM(5);\r
+\r
+ cv::gpu::setDevice(devInfo.deviceID());\r
+ }\r
+};\r
+\r
+TEST_P(ResizeArea, Accuracy)\r
+{\r
+ cv::Mat src = randomMat(size, type);\r
+\r
+ cv::gpu::GpuMat dst = createMat(cv::Size(cv::saturate_cast<int>(src.cols * coeff), cv::saturate_cast<int>(src.rows * coeff)), type, useRoi);\r
+ cv::gpu::resize(loadMat(src, useRoi), dst, cv::Size(), coeff, coeff, interpolation);\r
+\r
+ cv::Mat dst_cpu;\r
+ cv::resize(src, dst_cpu, cv::Size(), coeff, coeff, interpolation);\r
+\r
+ EXPECT_MAT_NEAR(dst_cpu, dst, src.depth() == CV_32F ? 1e-2 : 1.0);\r
+}\r
+\r
+INSTANTIATE_TEST_CASE_P(GPU_ImgProc, ResizeArea, testing::Combine(\r
+ ALL_DEVICES,\r
+ DIFFERENT_SIZES,\r
+ testing::Values(MatType(CV_8UC3), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)),\r
+ testing::Values(/*0.3,*/0.5),\r
+ testing::Values(Interpolation(cv::INTER_AREA)),\r
+ WHOLE_SUBMAT));\r
+\r
///////////////////////////////////////////////////////////////////\r
// Test NPP\r
\r
\r
CV_ENUM(NormCode, cv::NORM_INF, cv::NORM_L1, cv::NORM_L2, cv::NORM_TYPE_MASK, cv::NORM_RELATIVE, cv::NORM_MINMAX)\r
\r
-CV_ENUM(Interpolation, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC)\r
+CV_ENUM(Interpolation, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC, cv::INTER_AREA)\r
\r
CV_ENUM(BorderType, cv::BORDER_REFLECT101, cv::BORDER_REPLICATE, cv::BORDER_CONSTANT, cv::BORDER_REFLECT, cv::BORDER_WRAP)\r
#define ALL_BORDER_TYPES testing::Values(BorderType(cv::BORDER_REFLECT101), BorderType(cv::BORDER_REPLICATE), BorderType(cv::BORDER_CONSTANT), BorderType(cv::BORDER_REFLECT), BorderType(cv::BORDER_WRAP))\r
VecOp vecOp;
int x = vecOp((const uchar**)src, (uchar*)dst, (const uchar*)beta, width);
- #if CV_ENABLE_UNROLLED
- for( ; x <= width - 4; x += 4 )
+ #if CV_ENABLE_UNROLLED
+ for( ; x <= width - 4; x += 4 )
{
WT t0, t1;
t0 = S0[x]*b0 + S1[x]*b1;
CastOp castOp;
VecOp vecOp;
int k, x = vecOp((const uchar**)src, (uchar*)dst, (const uchar*)beta, width);
- #if CV_ENABLE_UNROLLED
+ #if CV_ENABLE_UNROLLED
for( ; x <= width - 4; x += 4 )
{
WT b = beta[0];
if( k0 < ksize )
hresize( srows + k0, rows + k0, ksize - k0, xofs, alpha,
ssize.width, dsize.width, cn, xmin, xmax );
- vresize( (const WT**)rows, (T*)(dst.data + dst.step*dy), beta, dsize.width );
+ vresize( (const WT**)rows, (T*)(dst.data + dst.step*dy), beta, dsize.width );
}
}
{
const T* S = (const T*)(src.data + src.step*sy0) + xofs[dx];
WT sum = 0;
- k=0;
- #if CV_ENABLE_UNROLLED
+ k=0;
+ #if CV_ENABLE_UNROLLED
for( ; k <= area - 4; k += 4 )
sum += S[ofs[k]] + S[ofs[k+1]] + S[ofs[k+2]] + S[ofs[k+3]];
#endif
resizeGeneric_<
HResizeLinear<uchar, int, short,
INTER_RESIZE_COEF_SCALE,
- HResizeLinearVec_8u32s>,
+ HResizeLinearVec_8u32s>,
VResizeLinear<uchar, int, short,
FixedPtCast<int, uchar, INTER_RESIZE_COEF_BITS*2>,
- VResizeLinearVec_32s8u> >,
- 0,
+ VResizeLinearVec_32s8u> >,
+ 0,
resizeGeneric_<
HResizeLinear<ushort, float, float, 1,
HResizeLinearVec_16u32f>,
HResizeLinearVec_16s32f>,
VResizeLinear<short, float, float, Cast<float, short>,
VResizeLinearVec_32f16s> >,
- 0,
+ 0,
resizeGeneric_<
HResizeLinear<float, float, float, 1,
HResizeLinearVec_32f>,
HResizeCubic<short, float, float>,
VResizeCubic<short, float, float, Cast<float, short>,
VResizeCubicVec_32f16s> >,
- 0,
+ 0,
resizeGeneric_<
HResizeCubic<float, float, float>,
VResizeCubic<float, float, float, Cast<float, float>,
resizeGeneric_<HResizeLanczos4<ushort, float, float>,
VResizeLanczos4<ushort, float, float, Cast<float, ushort>,
VResizeNoVec> >,
- resizeGeneric_<HResizeLanczos4<short, float, float>,
+ resizeGeneric_<HResizeLanczos4<short, float, float>,
VResizeLanczos4<short, float, float, Cast<float, short>,
VResizeNoVec> >,
- 0,
+ 0,
resizeGeneric_<HResizeLanczos4<float, float, float>,
VResizeLanczos4<float, float, float, Cast<float, float>,
VResizeNoVec> >,
static ResizeAreaFastFunc areafast_tab[] =
{
resizeAreaFast_<uchar, int>, 0,
- resizeAreaFast_<ushort, float>,
- resizeAreaFast_<short, float>,
+ resizeAreaFast_<ushort, float>,
+ resizeAreaFast_<short, float>,
0,
resizeAreaFast_<float, float>,
resizeAreaFast_<double, double>,
assert( k < ssize.width*2 );
xofs[k].di = dx*cn;
xofs[k].si = sx2*cn;
- xofs[k++].alpha = (float)((fsx2 - sx2)*scale);
+ xofs[k++].alpha = (float)(min(fsx2 - sx2, 1.)*scale);
}
}
if( !CV_ARE_TYPES_EQ( src, dst ))
CV_Error( CV_StsUnmatchedFormats, "" );
- ssize.width = src->cols;
+ ssize.width = src->cols;
ssize.height = src->rows;
dsize.width = dst->cols;
dsize.height = dst->rows;
ASSERT_EQ(line2.size(), (size_t)6);
}
+TEST(Imgproc_resize_area, regression)
+{
+ static ushort input_data[16 * 16] = {
+ 90, 94, 80, 3, 231, 2, 186, 245, 188, 165, 10, 19, 201, 169, 8, 228,
+ 86, 5, 203, 120, 136, 185, 24, 94, 81, 150, 163, 137, 88, 105, 132, 132,
+ 236, 48, 250, 218, 19, 52, 54, 221, 159, 112, 45, 11, 152, 153, 112, 134,
+ 78, 133, 136, 83, 65, 76, 82, 250, 9, 235, 148, 26, 236, 179, 200, 50,
+ 99, 51, 103, 142, 201, 65, 176, 33, 49, 226, 177, 109, 46, 21, 67, 130,
+ 54, 125, 107, 154, 145, 51, 199, 189, 161, 142, 231, 240, 139, 162, 240, 22,
+ 231, 86, 79, 106, 92, 47, 146, 156, 36, 207, 71, 33, 2, 244, 221, 71,
+ 44, 127, 71, 177, 75, 126, 68, 119, 200, 129, 191, 251, 6, 236, 247, 6,
+ 133, 175, 56, 239, 147, 221, 243, 154, 242, 82, 106, 99, 77, 158, 60, 229,
+ 2, 42, 24, 174, 27, 198, 14, 204, 246, 251, 141, 31, 114, 163, 29, 147,
+ 121, 53, 74, 31, 147, 189, 42, 98, 202, 17, 228, 123, 209, 40, 77, 49,
+ 112, 203, 30, 12, 205, 25, 19, 106, 145, 185, 163, 201, 237, 223, 247, 38,
+ 33, 105, 243, 117, 92, 179, 204, 248, 160, 90, 73, 126, 2, 41, 213, 204,
+ 6, 124, 195, 201, 230, 187, 210, 167, 48, 79, 123, 159, 145, 218, 105, 209,
+ 240, 152, 136, 235, 235, 164, 157, 9, 152, 38, 27, 209, 120, 77, 238, 196,
+ 240, 233, 10, 241, 90, 67, 12, 79, 0, 43, 58, 27, 83, 199, 190, 182};
+
+ static ushort expected_data = {
+ // row 1
+ 120, 100, 151, 101, 103,//?
+ // row 2
+ 106, 115, 141, 130, suturate_cast<ushort>(0),
+ // row 3
+ 91, 136, 170, 114, suturate_cast<ushort>(0),
+ // row 4
+ 104, 122, 131, 147, suturate_cast<ushort>(0),
+ // row 5
+ suturate_cast<ushort>(0),suturate_cast<ushort>(0),suturate_cast<ushort>(0),suturate_cast<ushort>(0),suturate_cast<ushort>(0)
+ };
+}
+
//////////////////////////////////////////////////////////////////////////
TEST(Imgproc_Resize, accuracy) { CV_ResizeTest test; test.safe_run(); }
projects / platform / upstream / opencv.git / commitdiff