From b2079d6de15982f6205546e3597653f42e2f5414 Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <no@email>
Date: Mon, 12 Sep 2011 10:13:14 +0000
Subject: [PATCH] implemented gpu::resize for all types

---
 modules/gpu/include/opencv2/gpu/gpu.hpp |   3 +-
 modules/gpu/perf/perf_imgproc.cpp       |   6 +-
 modules/gpu/src/cuda/imgproc.cu         | 216 ++++++++++++++++++++++++++++++++
 modules/gpu/src/imgproc.cpp             |  73 +++++++----
 modules/gpu/test/test_imgproc.cpp       |  12 +-
 samples/gpu/performance/tests.cpp       | 128 ++++++++++++++++++-
 6 files changed, 401 insertions(+), 37 deletions(-)

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 250c630..3c7dcc7 100644
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -633,8 +633,7 @@ namespace cv
         CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh, double maxval, int type, Stream& stream = Stream::Null());
 
         //! resizes the image
-        //! Supports INTER_NEAREST, INTER_LINEAR
-        //! supports CV_8UC1, CV_8UC4 types
+        //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC
         CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null());
 
         //! warps the image using affine transformation
diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp
index 67e7466..f4d63d6 100644
--- a/modules/gpu/perf/perf_imgproc.cpp
+++ b/modules/gpu/perf/perf_imgproc.cpp
@@ -245,8 +245,8 @@ PERF_TEST_P(DevInfo_Size_MatType, threshold, testing::Combine(testing::ValuesIn(
 
 PERF_TEST_P(DevInfo_Size_MatType_Interpolation_SizeCoeff, resize, testing::Combine(testing::ValuesIn(devices()),
                                                                                    testing::Values(GPU_TYPICAL_MAT_SIZES), 
-                                                                                   testing::Values(CV_8UC1, CV_8UC4),
-                                                                                   testing::Values((int)INTER_NEAREST, (int)INTER_LINEAR),
+                                                                                   testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                                                                                   testing::Values((int)INTER_NEAREST, (int)INTER_LINEAR, (int)INTER_CUBIC),
                                                                                    testing::Values(0.5, 2.0)))
 {
     DeviceInfo devInfo = std::tr1::get<0>(GetParam());
@@ -264,7 +264,7 @@ PERF_TEST_P(DevInfo_Size_MatType_Interpolation_SizeCoeff, resize, testing::Combi
     GpuMat src(src_host);
     GpuMat dst;
 
-    declare.time(0.5).iterations(100);
+    declare.time(1.0).iterations(100);
 
     SIMPLE_TEST_CYCLE()
     {
diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 44a54ca..d76f93b 100644
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -252,6 +252,222 @@ namespace cv { namespace gpu { namespace imgproc
     template void remap_gpu<float3>(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2D& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream);
     template void remap_gpu<float4>(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2D& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream);
 
+/////////////////////////////////// Resize ///////////////////////////////////////////////
+
+    template <typename Ptr2D, typename T> __global__ void resize(const Ptr2D src, float fx, float fy, DevMem2D_<T> dst)
+    {
+        const int x = blockDim.x * blockIdx.x + threadIdx.x;
+        const int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+        if (x < dst.cols && y < dst.rows)
+        {
+            const float xcoo = x / fx;
+            const float ycoo = y / fy;
+
+            dst.ptr(y)[x] = saturate_cast<T>(src(ycoo, xcoo));
+        }
+    }
+    template <typename Ptr2D, typename T> __global__ void resizeNN(const Ptr2D src, float fx, float fy, DevMem2D_<T> dst)
+    {
+        const int x = blockDim.x * blockIdx.x + threadIdx.x;
+        const int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+        if (x < dst.cols && y < dst.rows)
+        {
+            const float xcoo = x / fx;
+            const float ycoo = y / fy;
+
+            dst.ptr(y)[x] = src(__float2int_rd(ycoo), __float2int_rd(xcoo));
+        }
+    }
+
+    template <template <typename> class Filter, typename T> struct ResizeDispatcherStream
+    {
+        static void call(const DevMem2D_<T>& src, float fx, float fy, const DevMem2D_<T>& dst, cudaStream_t stream)
+        {            
+            dim3 block(32, 8);
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
+
+            BrdReplicate<T> brd(src.rows, src.cols);
+            BorderReader< PtrStep_<T>, BrdReplicate<T> > brdSrc(src, brd);
+            Filter< BorderReader< PtrStep_<T>, BrdReplicate<T> > > filter_src(brdSrc);
+
+            resize<<<grid, block, 0, stream>>>(filter_src, fx, fy, dst);
+            cudaSafeCall( cudaGetLastError() );
+        }
+    };
+    template <typename T> struct ResizeDispatcherStream<PointFilter, T>
+    {
+        static void call(const DevMem2D_<T>& src, float fx, float fy, const DevMem2D_<T>& dst, cudaStream_t stream)
+        {            
+            dim3 block(32, 8);
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
+
+            BrdReplicate<T> brd(src.rows, src.cols);
+            BorderReader< PtrStep_<T>, BrdReplicate<T> > brdSrc(src, brd);
+
+            resizeNN<<<grid, block, 0, stream>>>(brdSrc, fx, fy, dst);
+            cudaSafeCall( cudaGetLastError() );
+        }
+    };
+    
+    template <template <typename> class Filter, typename T> struct ResizeDispatcherNonStream
+    {
+        static void call(const DevMem2D_<T>& src, float fx, float fy, const DevMem2D_<T>& dst)
+        {            
+            dim3 block(32, 8);
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
+
+            BrdReplicate<T> brd(src.rows, src.cols);
+            BorderReader< PtrStep_<T>, BrdReplicate<T> > brdSrc(src, brd);
+            Filter< BorderReader< PtrStep_<T>, BrdReplicate<T> > > filter_src(brdSrc);
+
+            resize<<<grid, block>>>(filter_src, fx, fy, dst);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+        }
+    };
+    template <typename T> struct ResizeDispatcherNonStream<PointFilter, T>
+    {
+        static void call(const DevMem2D_<T>& src, float fx, float fy, const DevMem2D_<T>& dst)
+        {            
+            dim3 block(32, 8);
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
+
+            BrdReplicate<T> brd(src.rows, src.cols);
+            BorderReader< PtrStep_<T>, BrdReplicate<T> > brdSrc(src, brd);
+
+            resizeNN<<<grid, block>>>(brdSrc, fx, fy, dst);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+        }
+    };
+
+#define OPENCV_GPU_IMPLEMENT_RESIZE_TEX(type) \
+    texture< type , cudaTextureType2D> tex_resize_ ## type (0, cudaFilterModePoint, cudaAddressModeClamp); \
+    struct tex_resize_ ## type ## _reader \
+    { \
+        typedef type elem_type; \
+        typedef int index_type; \
+        __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \
+        { \
+            return tex2D(tex_resize_ ## type , x, y); \
+        } \
+    }; \
+    template <template <typename> class Filter> struct ResizeDispatcherNonStream<Filter, type> \
+    { \
+        static void call(const DevMem2D_< type >& src, float fx, float fy, const DevMem2D_< type >& dst) \
+        { \
+            dim3 block(32, 8); \
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \
+            TextureBinder texHandler(&tex_resize_ ## type , src); \
+            tex_resize_ ## type ##_reader texSrc; \
+            Filter< tex_resize_ ## type ##_reader > filter_src(texSrc); \
+            resize<<<grid, block>>>(filter_src, fx, fy, dst); \
+            cudaSafeCall( cudaGetLastError() ); \
+            cudaSafeCall( cudaDeviceSynchronize() ); \
+        } \
+    }; \
+    template <> struct ResizeDispatcherNonStream<PointFilter, type> \
+    { \
+        static void call(const DevMem2D_< type >& src, float fx, float fy, const DevMem2D_< type >& dst) \
+        { \
+            dim3 block(32, 8); \
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \
+            TextureBinder texHandler(&tex_resize_ ## type , src); \
+            tex_resize_ ## type ##_reader texSrc; \
+            resizeNN<<<grid, block>>>(texSrc, fx, fy, dst); \
+            cudaSafeCall( cudaGetLastError() ); \
+            cudaSafeCall( cudaDeviceSynchronize() ); \
+        } \
+    };
+    
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(uchar)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(uchar2)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(uchar4)
+    
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(schar)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(char2)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(char4)
+    
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(ushort)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(ushort2)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(ushort4)
+    
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(short)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(short2)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(short4)
+    
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(int)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(int2)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(int4)
+    
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(float)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(float2)
+    OPENCV_GPU_IMPLEMENT_RESIZE_TEX(float4)
+    
+#undef OPENCV_GPU_IMPLEMENT_RESIZE_TEX
+
+    template <template <typename> class Filter, typename T> struct ResizeDispatcher
+    { 
+        static void call(const DevMem2D_<T>& src, float fx, float fy, const DevMem2D_<T>& dst, cudaStream_t stream)
+        {
+            if (stream == 0)
+                ResizeDispatcherNonStream<Filter, T>::call(src, fx, fy, dst);
+            else
+                ResizeDispatcherStream<Filter, T>::call(src, fx, fy, dst, stream);
+        }
+    };
+
+    template <typename T> void resize_gpu(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream)
+    {
+        typedef void (*caller_t)(const DevMem2D_<T>& src, float fx, float fy, const DevMem2D_<T>& dst, cudaStream_t stream);
+
+        static const caller_t callers[3] = 
+        {
+            ResizeDispatcher<PointFilter, T>::call, ResizeDispatcher<LinearFilter, T>::call, ResizeDispatcher<CubicFilter, T>::call
+        };
+
+        callers[interpolation](static_cast< DevMem2D_<T> >(src), fx, fy, static_cast< DevMem2D_<T> >(dst), stream);
+    }
+
+    template void resize_gpu<uchar >(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<uchar2>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<uchar3>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<uchar4>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    
+    template void resize_gpu<schar>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<char2>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<char3>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<char4>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    
+    template void resize_gpu<ushort >(const DevMem2D& src,float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<ushort2>(const DevMem2D& src,float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<ushort3>(const DevMem2D& src,float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<ushort4>(const DevMem2D& src,float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    
+    template void resize_gpu<short >(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<short2>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<short3>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<short4>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    
+    template void resize_gpu<uint >(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<uint2>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<uint3>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<uint4>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    
+    template void resize_gpu<int >(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<int2>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<int3>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<int4>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    
+    template void resize_gpu<float >(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<float2>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<float3>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    template void resize_gpu<float4>(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+
 /////////////////////////////////// MeanShiftfiltering ///////////////////////////////////////////////
 
     texture<uchar4, 2> tex_meanshift;
diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp
index 19ac1a1..b9d3607 100644
--- a/modules/gpu/src/imgproc.cpp
+++ b/modules/gpu/src/imgproc.cpp
@@ -272,14 +272,28 @@ void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q,
 ////////////////////////////////////////////////////////////////////////
 // resize
 
+namespace cv { namespace gpu {  namespace imgproc
+{
+    template <typename T> void resize_gpu(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+}}}
+
 void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation, Stream& s)
 {
-    static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR/*, NPPI_INTER_CUBIC, 0, NPPI_INTER_LANCZOS*/};
+    using namespace cv::gpu::imgproc;
 
-    CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);
-    CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR/* || interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4*/);
+    typedef void (*caller_t)(const DevMem2D& src, float fx, float fy, const DevMem2D& dst, int interpolation, cudaStream_t stream);
+    static const caller_t callers[6][4] = 
+    {
+        {resize_gpu<uchar>, resize_gpu<uchar2>, resize_gpu<uchar3>, resize_gpu<uchar4>},
+        {resize_gpu<schar>, resize_gpu<char2>, resize_gpu<char3>, resize_gpu<char4>},
+        {resize_gpu<ushort>, resize_gpu<ushort2>, resize_gpu<ushort3>, resize_gpu<ushort4>},
+        {resize_gpu<short>, resize_gpu<short2>, resize_gpu<short3>, resize_gpu<short4>},
+        {resize_gpu<int>, resize_gpu<int2>, resize_gpu<int3>, resize_gpu<int4>},
+        {resize_gpu<float>, resize_gpu<float2>, resize_gpu<float3>, resize_gpu<float4>}
+    };
 
-    CV_Assert( src.size().area() > 0 );
+    CV_Assert( src.depth() <= CV_32F && src.channels() <= 4 );
+    CV_Assert( interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC );
     CV_Assert( !(dsize == Size()) || (fx > 0 && fy > 0) );
 
     if( dsize == Size() )
@@ -294,34 +308,43 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub
 
     dst.create(dsize, src.type());
 
-    NppiSize srcsz;
-    srcsz.width  = src.cols;
-    srcsz.height = src.rows;
-    NppiRect srcrect;
-    srcrect.x = srcrect.y = 0;
-    srcrect.width  = src.cols;
-    srcrect.height = src.rows;
-    NppiSize dstsz;
-    dstsz.width  = dst.cols;
-    dstsz.height = dst.rows;
-
     cudaStream_t stream = StreamAccessor::getStream(s);
 
-    NppStreamHandler h(stream);
-
-    if (src.type() == CV_8UC1)
+    if ((src.type() == CV_8UC1 || src.type() == CV_8UC4) && (interpolation == INTER_NEAREST || interpolation == INTER_LINEAR))
     {
-        nppSafeCall( nppiResize_8u_C1R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcrect,
-            dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, fx, fy, npp_inter[interpolation]) );
+        static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC, 0, NPPI_INTER_LANCZOS};
+
+        NppiSize srcsz;
+        srcsz.width  = src.cols;
+        srcsz.height = src.rows;
+        NppiRect srcrect;
+        srcrect.x = srcrect.y = 0;
+        srcrect.width  = src.cols;
+        srcrect.height = src.rows;
+        NppiSize dstsz;
+        dstsz.width  = dst.cols;
+        dstsz.height = dst.rows;
+
+        NppStreamHandler h(stream);
+
+        if (src.type() == CV_8UC1)
+        {
+            nppSafeCall( nppiResize_8u_C1R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcrect,
+                dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, fx, fy, npp_inter[interpolation]) );
+        }
+        else
+        {
+            nppSafeCall( nppiResize_8u_C4R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcrect,
+                dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, fx, fy, npp_inter[interpolation]) );
+        }
+
+        if (stream == 0)
+            cudaSafeCall( cudaDeviceSynchronize() );
     }
     else
     {
-        nppSafeCall( nppiResize_8u_C4R(src.ptr<Npp8u>(), srcsz, static_cast<int>(src.step), srcrect,
-            dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, fx, fy, npp_inter[interpolation]) );
+        callers[src.depth()][src.channels() - 1](src, static_cast<float>(fx), static_cast<float>(fy), dst, interpolation, stream);
     }
-
-    if (stream == 0)
-        cudaSafeCall( cudaDeviceSynchronize() );
 }
 
 ////////////////////////////////////////////////////////////////////////
diff --git a/modules/gpu/test/test_imgproc.cpp b/modules/gpu/test/test_imgproc.cpp
index 094ed9a..f476fbb 100644
--- a/modules/gpu/test/test_imgproc.cpp
+++ b/modules/gpu/test/test_imgproc.cpp
@@ -137,7 +137,7 @@ struct Resize : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int
 
         size = cv::Size(rng.uniform(20, 150), rng.uniform(20, 150));
 
-        src = cvtest::randomMat(rng, size, type, 0.0, 127.0, false);
+        src = cvtest::randomMat(rng, size, type, 0.0, CV_MAT_DEPTH(type) == CV_32F ? 1.0 : 255.0, false);
 
         cv::resize(src, dst_gold1, cv::Size(), 2.0, 2.0, interpolation);
         cv::resize(src, dst_gold2, cv::Size(), 0.5, 0.5, interpolation);
@@ -146,7 +146,7 @@ struct Resize : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int
 
 TEST_P(Resize, Accuracy)
 {
-    static const char* interpolations[] = {"INTER_NEAREST", "INTER_LINEAR"};
+    static const char* interpolations[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC"};
     const char* interpolationStr = interpolations[interpolation];
 
     PRINT_PARAM(devInfo);
@@ -169,14 +169,14 @@ TEST_P(Resize, Accuracy)
         gpuRes2.download(dst2);
     );
 
-    EXPECT_MAT_SIMILAR(dst_gold1, dst1, 0.5);
-    EXPECT_MAT_SIMILAR(dst_gold2, dst2, 0.5);
+    EXPECT_MAT_SIMILAR(dst_gold1, dst1, 0.2);
+    EXPECT_MAT_SIMILAR(dst_gold2, dst2, 0.2);
 }
 
 INSTANTIATE_TEST_CASE_P(ImgProc, Resize, testing::Combine(
                         testing::ValuesIn(devices()), 
-                        testing::Values(CV_8UC1, CV_8UC4), 
-                        testing::Values((int)cv::INTER_NEAREST, (int)cv::INTER_LINEAR)));
+                        testing::Values(CV_8UC1, CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), 
+                        testing::Values((int)cv::INTER_NEAREST, (int)cv::INTER_LINEAR, (int)cv::INTER_CUBIC)));
 
 ///////////////////////////////////////////////////////////////////////////////////////////////////////
 // remap
diff --git a/samples/gpu/performance/tests.cpp b/samples/gpu/performance/tests.cpp
index 5f70236..813f527 100644
--- a/samples/gpu/performance/tests.cpp
+++ b/samples/gpu/performance/tests.cpp
@@ -592,9 +592,81 @@ TEST(resize)
 
     for (int size = 1000; size <= 3000; size += 1000)
     {
-        SUBTEST << "size " << size;
+        SUBTEST << "size " << size << ", 8UC1, up";
+
+        gen(src, size, size, CV_8U, 0, 256);
+        dst.create(size * 2, size * 2, CV_8U);
+
+        CPU_ON;
+        resize(src, dst, dst.size());
+        CPU_OFF;
+
+        d_src = src;
+        d_dst.create(size * 2, size * 2, CV_8U);
+
+        GPU_ON;
+        gpu::resize(d_src, d_dst, d_dst.size());
+        GPU_OFF;
+    }
+    for (int size = 1000; size <= 3000; size += 1000)
+    {
+        SUBTEST << "size " << size << ", 8UC1, down";
 
         gen(src, size, size, CV_8U, 0, 256);
+        dst.create(size / 2, size / 2, CV_8U);
+
+        CPU_ON;
+        resize(src, dst, dst.size());
+        CPU_OFF;
+
+        d_src = src;
+        d_dst.create(size / 2, size / 2, CV_8U);
+
+        GPU_ON;
+        gpu::resize(d_src, d_dst, d_dst.size());
+        GPU_OFF;
+    }
+    for (int size = 1000; size <= 3000; size += 1000)
+    {
+        SUBTEST << "size " << size << ", 8UC3, up";
+
+        gen(src, size, size, CV_8UC3, 0, 256);
+        dst.create(size * 2, size * 2, CV_8U);
+
+        CPU_ON;
+        resize(src, dst, dst.size());
+        CPU_OFF;
+
+        d_src = src;
+        d_dst.create(size * 2, size * 2, CV_8U);
+
+        GPU_ON;
+        gpu::resize(d_src, d_dst, d_dst.size());
+        GPU_OFF;
+    }
+    for (int size = 1000; size <= 3000; size += 1000)
+    {
+        SUBTEST << "size " << size << ", 8UC3, down";
+
+        gen(src, size, size, CV_8UC3, 0, 256);
+        dst.create(size / 2, size / 2, CV_8U);
+
+        CPU_ON;
+        resize(src, dst, dst.size());
+        CPU_OFF;
+
+        d_src = src;
+        d_dst.create(size / 2, size / 2, CV_8U);
+
+        GPU_ON;
+        gpu::resize(d_src, d_dst, d_dst.size());
+        GPU_OFF;
+    }
+    for (int size = 1000; size <= 3000; size += 1000)
+    {
+        SUBTEST << "size " << size << ", 8UC4, up";
+
+        gen(src, size, size, CV_8UC4, 0, 256);
         dst.create(size * 2, size * 2, CV_8U);
 
         CPU_ON;
@@ -608,6 +680,60 @@ TEST(resize)
         gpu::resize(d_src, d_dst, d_dst.size());
         GPU_OFF;
     }
+    for (int size = 1000; size <= 3000; size += 1000)
+    {
+        SUBTEST << "size " << size << ", 8UC4, down";
+
+        gen(src, size, size, CV_8UC4, 0, 256);
+        dst.create(size / 2, size / 2, CV_8U);
+
+        CPU_ON;
+        resize(src, dst, dst.size());
+        CPU_OFF;
+
+        d_src = src;
+        d_dst.create(size / 2, size / 2, CV_8U);
+
+        GPU_ON;
+        gpu::resize(d_src, d_dst, d_dst.size());
+        GPU_OFF;
+    }
+    for (int size = 1000; size <= 3000; size += 1000)
+    {
+        SUBTEST << "size " << size << ", 32FC1, up";
+
+        gen(src, size, size, CV_32FC1, 0, 256);
+        dst.create(size * 2, size * 2, CV_8U);
+
+        CPU_ON;
+        resize(src, dst, dst.size());
+        CPU_OFF;
+
+        d_src = src;
+        d_dst.create(size * 2, size * 2, CV_8U);
+
+        GPU_ON;
+        gpu::resize(d_src, d_dst, d_dst.size());
+        GPU_OFF;
+    }
+    for (int size = 1000; size <= 3000; size += 1000)
+    {
+        SUBTEST << "size " << size << ", 32FC1, down";
+
+        gen(src, size, size, CV_32FC1, 0, 256);
+        dst.create(size / 2, size / 2, CV_8U);
+
+        CPU_ON;
+        resize(src, dst, dst.size());
+        CPU_OFF;
+
+        d_src = src;
+        d_dst.create(size / 2, size / 2, CV_8U);
+
+        GPU_ON;
+        gpu::resize(d_src, d_dst, d_dst.size());
+        GPU_OFF;
+    }
 }
 
 
-- 
2.7.4