             { \
                 static void call(PtrStepSz< type > src, PtrStepSz< type > srcWhole, int xoff, int yoff, float fx, float fy, PtrStepSz< type > dst) \
                 { \
-                    dim3 block(32, 8); \
-                    dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \
+                    const dim3 block(32, 8); \
+                    const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \
                     bindTexture(&tex_resize_ ## type, srcWhole); \
-                    tex_resize_ ## type ## _reader texSrc(xoff, yoff); \
+                    tex_resize_ ## type ## _reader texSrc; \
+                    texSrc.xoff = xoff; \
+                    texSrc.yoff = yoff; \
                     if (srcWhole.cols == src.cols && srcWhole.rows == src.rows) \
                     { \
                         Filter<tex_resize_ ## type ## _reader> filteredSrc(texSrc); \

            { \
                typedef type elem_type; \
                typedef int index_type; \
                const int xoff; \
                const int yoff; \
                __host__ tex_resize_ ## type ## _reader(int xoff_, int yoff_) : xoff(xoff_), yoff(yoff_) {} \
                __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \
                { \
                    return tex2D(tex_resize_ ## type, x + xoff, y + yoff); \

            }
        };

        template <typename Ptr2D, typename T> __global__ void resize_area(const Ptr2D src, 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)
            {
                dst(y, x) = saturate_cast<T>(src(y, x));
            }
        }

        template <typename T> struct ResizeAreaDispatcher
        {
            static void call(PtrStepSz<T> src, PtrStepSz<T>, int, int, float fx, float fy, PtrStepSz<T> dst, cudaStream_t stream)
            {
                const int iscale_x = (int) round(fx);
                const int iscale_y = (int) round(fy);

                const dim3 block(32, 8);
                const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));

                if (std::abs(fx - iscale_x) < FLT_MIN && std::abs(fy - iscale_y) < FLT_MIN)
                {
                    BrdConstant<T> brd(src.rows, src.cols);
                    BorderReader< PtrStep<T>, BrdConstant<T> > brdSrc(src, brd);
                    IntegerAreaFilter< BorderReader< PtrStep<T>, BrdConstant<T> > > filteredSrc(brdSrc, fx, fy);
                    resize_area<<<grid, block, 0, stream>>>(filteredSrc, dst);
                }
                else
                {
                    BrdConstant<T> brd(src.rows, src.cols);
                    BorderReader< PtrStep<T>, BrdConstant<T> > brdSrc(src, brd);
                    AreaFilter< BorderReader< PtrStep<T>, BrdConstant<T> > > filteredSrc(brdSrc, fx, fy);
                    resize_area<<<grid, block, 0, stream>>>(filteredSrc, dst);
                }

                cudaSafeCall( cudaGetLastError() );
                if (stream == 0)
                    cudaSafeCall( cudaDeviceSynchronize() );
            }
        };

