 {
     namespace imgproc
     {
+        template <typename T> __global__ void resize_nearest(const PtrStep<T> src, const float fx, const float fy, PtrStepSz<T> dst)
+        {
+            const int dst_x = blockDim.x * blockIdx.x + threadIdx.x;
+            const int dst_y = blockDim.y * blockIdx.y + threadIdx.y;
+
+            if (dst_x < dst.cols && dst_y < dst.rows)
+            {
+                const float src_x = dst_x * fx;
+                const float src_y = dst_y * fy;
+
+                dst(dst_y, dst_x) = src(__float2int_rz(src_y), __float2int_rz(src_x));
+            }
+        }
+
+        template <typename T> __global__ void resize_linear(const PtrStepSz<T> src, const float fx, const float fy, PtrStepSz<T> dst)
+        {
+            typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type;
+
+            const int dst_x = blockDim.x * blockIdx.x + threadIdx.x;
+            const int dst_y = blockDim.y * blockIdx.y + threadIdx.y;
+
+            if (dst_x < dst.cols && dst_y < dst.rows)
+            {
+                const float src_x = dst_x * fx;
+                const float src_y = dst_y * fy;
+
+                work_type out = VecTraits<work_type>::all(0);
+
+                const int x1 = __float2int_rd(src_x);
+                const int y1 = __float2int_rd(src_y);
+                const int x2 = x1 + 1;
+                const int y2 = y1 + 1;
+                const int x2_read = ::min(x2, src.cols - 1);
+                const int y2_read = ::min(y2, src.rows - 1);
+
+                T src_reg = src(y1, x1);
+                out = out + src_reg * ((x2 - src_x) * (y2 - src_y));
+
+                src_reg = src(y1, x2_read);
+                out = out + src_reg * ((src_x - x1) * (y2 - src_y));
+
+                src_reg = src(y2_read, x1);
+                out = out + src_reg * ((x2 - src_x) * (src_y - y1));
+
+                src_reg = src(y2_read, x2_read);
+                out = out + src_reg * ((src_x - x1) * (src_y - y1));
+
+                dst(dst_y, dst_x) = saturate_cast<T>(out);
+            }
+        }
+
         template <class Ptr2D, typename T> __global__ void resize(const Ptr2D src, const float fx, const float fy, PtrStepSz<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(y, x) = saturate_cast<T>(src(ycoo, xcoo));
+            const int dst_x = blockDim.x * blockIdx.x + threadIdx.x;
+            const int dst_y = blockDim.y * blockIdx.y + threadIdx.y;
+
+            if (dst_x < dst.cols && dst_y < dst.rows)
+            {
+                const float src_x = dst_x * fx;
+                const float src_y = dst_y * fy;
+
+                dst(dst_y, dst_x) = src(src_y, src_x);
             }
         }
 
         template <template <typename> class Filter, typename T> struct ResizeDispatcherStream
         {
             static void call(PtrStepSz<T> src, float fx, float fy, PtrStepSz<T> dst, cudaStream_t stream)
-            {
-                const dim3 block(32, 8);
-                const 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> > > filteredSrc(brdSrc, fx, fy);
-
-                resize<<<grid, block, 0, stream>>>(filteredSrc, fx, fy, dst);
-                cudaSafeCall( cudaGetLastError() );
-            }
-        };
-
-        template <template <typename> class Filter, typename T> struct ResizeDispatcherNonStream
-        {
-            static void call(PtrStepSz<T> src, PtrStepSz<T>, int, int, float fx, float fy, PtrStepSz<T> dst)
             {
                 const dim3 block(32, 8);
                 const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));

                ResizeDispatcher<CubicFilter, T>::call,
                ResizeAreaDispatcher<T>::call
            };
            // chenge to linear if area interpolation upscaling
            if (interpolation == 3 && (fx <= 1.f || fy <= 1.f))
                interpolation = 1;


                BorderReader< PtrStep<T>, BrdReplicate<T> > brdSrc(src, brd);
                Filter< BorderReader< PtrStep<T>, BrdReplicate<T> > > filteredSrc(brdSrc);

                resize<<<grid, block, 0, stream>>>(filteredSrc, fx, fy, dst);
                cudaSafeCall( cudaGetLastError() );
            }
        };
        template <typename T> struct ResizeDispatcherStream<PointFilter, T>
        {
            static void call(PtrStepSz<T> src, float fx, float fy, PtrStepSz<T> dst, cudaStream_t stream)
            {
                const dim3 block(32, 8);
                const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));

                resize_nearest<<<grid, block, 0, stream>>>(src, fx, fy, dst);
                cudaSafeCall( cudaGetLastError() );
            }
        };
        template <typename T> struct ResizeDispatcherStream<LinearFilter, T>
        {
            static void call(PtrStepSz<T> src, float fx, float fy, PtrStepSz<T> dst, cudaStream_t stream)
            {
                const dim3 block(32, 8);
                const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));

                resize_linear<<<grid, block, 0, stream>>>(src, fx, fy, dst);
                cudaSafeCall( cudaGetLastError() );
            }
        };

        template <template <typename> class Filter, typename T> struct ResizeDispatcherNonStream
        {
            static void call(PtrStepSz<T> src, PtrStepSz<T>, int, int, float fx, float fy, PtrStepSz<T> dst)
            {
                const dim3 block(32, 8);
                const 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> > > filteredSrc(brdSrc);

                resize<<<grid, block>>>(filteredSrc, fx, fy, dst);
                cudaSafeCall( cudaGetLastError() );

                cudaSafeCall( cudaDeviceSynchronize() );
            }
        };
        template <typename T> struct ResizeDispatcherNonStream<PointFilter, T>
        {
        static void call(PtrStepSz<T> src, PtrStepSz<T>, int, int, float fx, float fy, PtrStepSz<T> dst)
            {
                const dim3 block(32, 8);
                const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));

                resize_nearest<<<grid, block>>>(src, fx, fy, dst);
                cudaSafeCall( cudaGetLastError() );

                cudaSafeCall( cudaDeviceSynchronize() );
            }
        };
        template <typename T> struct ResizeDispatcherNonStream<LinearFilter, T>
        {
        static void call(PtrStepSz<T> src, PtrStepSz<T>, int, int, float fx, float fy, PtrStepSz<T> dst)
            {
                const dim3 block(32, 8);
                const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));

                resize_linear<<<grid, block>>>(src, fx, fy, dst);
                cudaSafeCall( cudaGetLastError() );

                cudaSafeCall( cudaDeviceSynchronize() );
