/*M/////////////////////////////////////////////////////////////////////////////////////// // // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. // // By downloading, copying, installing or using the software you agree to this license. // If you do not agree to this license, do not download, install, // copy or use the software. // // // License Agreement // For Open Source Computer Vision Library // // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. // Copyright (C) 2009, Willow Garage Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // Redistribution and use in source and binary forms, with or without modification, // are permitted provided that the following conditions are met: // // * Redistribution's of source code must retain the above copyright notice, // this list of conditions and the following disclaimer. // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation // and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. // // This software is provided by the copyright holders and contributors "as is" and // any express or implied warranties, including, but not limited to, the implied // warranties of merchantability and fitness for a particular purpose are disclaimed. // In no event shall the Intel Corporation or contributors be liable for any direct, // indirect, incidental, special, exemplary, or consequential damages // (including, but not limited to, procurement of substitute goods or services; // loss of use, data, or profits; or business interruption) however caused // and on any theory of liability, whether in contract, strict liability, // or tort (including negligence or otherwise) arising in any way out of // the use of this software, even if advised of the possibility of such damage. // //M*/ #if !defined CUDA_DISABLER #include #include "opencv2/gpu/device/common.hpp" #include "opencv2/gpu/device/border_interpolate.hpp" #include "opencv2/gpu/device/vec_traits.hpp" #include "opencv2/gpu/device/vec_math.hpp" #include "opencv2/gpu/device/saturate_cast.hpp" #include "opencv2/gpu/device/filters.hpp" namespace cv { namespace gpu { namespace device { // kernels template __global__ void resize_nearest(const PtrStep src, PtrStepSz dst, const float fy, const float fx) { 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 __global__ void resize_linear(const PtrStepSz src, PtrStepSz dst, const float fy, const float fx) { typedef typename TypeVec::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::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(out); } } template __global__ void resize(const Ptr2D src, PtrStepSz dst, const float fy, const float fx) { 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 __global__ void resize_area(const Ptr2D src, PtrStepSz 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) = src(y, x); } } // textures template struct TextureAccessor; #define OPENCV_GPU_IMPLEMENT_RESIZE_TEX(type) \ texture tex_resize_##type (0, cudaFilterModePoint, cudaAddressModeClamp); \ template <> struct TextureAccessor \ { \ typedef type elem_type; \ typedef int index_type; \ int xoff; \ int yoff; \ __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \ { \ return tex2D(tex_resize_##type, x + xoff, y + yoff); \ } \ __host__ static void bind(const PtrStepSz& mat) \ { \ bindTexture(&tex_resize_##type, mat); \ } \ }; OPENCV_GPU_IMPLEMENT_RESIZE_TEX(uchar) OPENCV_GPU_IMPLEMENT_RESIZE_TEX(uchar4) OPENCV_GPU_IMPLEMENT_RESIZE_TEX(ushort) OPENCV_GPU_IMPLEMENT_RESIZE_TEX(ushort4) OPENCV_GPU_IMPLEMENT_RESIZE_TEX(short) OPENCV_GPU_IMPLEMENT_RESIZE_TEX(short4) OPENCV_GPU_IMPLEMENT_RESIZE_TEX(float) OPENCV_GPU_IMPLEMENT_RESIZE_TEX(float4) #undef OPENCV_GPU_IMPLEMENT_RESIZE_TEX template TextureAccessor texAccessor(const PtrStepSz& mat, int yoff, int xoff) { TextureAccessor::bind(mat); TextureAccessor t; t.xoff = xoff; t.yoff = yoff; return t; } // callers for nearest interpolation template void call_resize_nearest_glob(const PtrStepSz& src, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { const dim3 block(32, 8); const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); resize_nearest<<>>(src, dst, fy, fx); cudaSafeCall( cudaGetLastError() ); if (stream == 0) cudaSafeCall( cudaDeviceSynchronize() ); } template void call_resize_nearest_tex(const PtrStepSz& /*src*/, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx) { const dim3 block(32, 8); const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); resize<<>>(texAccessor(srcWhole, yoff, xoff), dst, fy, fx); cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); } // callers for linear interpolation template void call_resize_linear_glob(const PtrStepSz& src, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { const dim3 block(32, 8); const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); resize_linear<<>>(src, dst, fy, fx); cudaSafeCall( cudaGetLastError() ); if (stream == 0) cudaSafeCall( cudaDeviceSynchronize() ); } template void call_resize_linear_tex(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx) { const dim3 block(32, 8); const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); if (srcWhole.data == src.data) { TextureAccessor texSrc = texAccessor(src, 0, 0); LinearFilter< TextureAccessor > filteredSrc(texSrc); resize<<>>(filteredSrc, dst, fy, fx); } else { TextureAccessor texSrc = texAccessor(srcWhole, yoff, xoff); BrdReplicate brd(src.rows, src.cols); BorderReader, BrdReplicate > brdSrc(texSrc, brd); LinearFilter< BorderReader, BrdReplicate > > filteredSrc(brdSrc); resize<<>>(filteredSrc, dst, fy, fx); } cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); } // callers for cubic interpolation template void call_resize_cubic_glob(const PtrStepSz& src, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { const dim3 block(32, 8); const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); BrdReplicate brd(src.rows, src.cols); BorderReader< PtrStep, BrdReplicate > brdSrc(src, brd); CubicFilter< BorderReader< PtrStep, BrdReplicate > > filteredSrc(brdSrc); resize<<>>(filteredSrc, dst, fy, fx); cudaSafeCall( cudaGetLastError() ); if (stream == 0) cudaSafeCall( cudaDeviceSynchronize() ); } template void call_resize_cubic_tex(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx) { const dim3 block(32, 8); const dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); if (srcWhole.data == src.data) { TextureAccessor texSrc = texAccessor(src, 0, 0); CubicFilter< TextureAccessor > filteredSrc(texSrc); resize<<>>(filteredSrc, dst, fy, fx); } else { TextureAccessor texSrc = texAccessor(srcWhole, yoff, xoff); BrdReplicate brd(src.rows, src.cols); BorderReader, BrdReplicate > brdSrc(texSrc, brd); CubicFilter< BorderReader, BrdReplicate > > filteredSrc(brdSrc); resize<<>>(filteredSrc, dst, fy, fx); } cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); } // ResizeNearestDispatcher template struct ResizeNearestDispatcher { static void call(const PtrStepSz& src, const PtrStepSz& /*srcWhole*/, int /*yoff*/, int /*xoff*/, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { call_resize_nearest_glob(src, dst, fy, fx, stream); } }; template struct SelectImplForNearest { static void call(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { if (stream) call_resize_nearest_glob(src, dst, fy, fx, stream); else { if (fx > 1 || fy > 1) call_resize_nearest_glob(src, dst, fy, fx, 0); else call_resize_nearest_tex(src, srcWhole, yoff, xoff, dst, fy, fx); } } }; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; template <> struct ResizeNearestDispatcher : SelectImplForNearest {}; // ResizeLinearDispatcher template struct ResizeLinearDispatcher { static void call(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { call_resize_linear_glob(src, dst, fy, fx, stream); } }; template struct SelectImplForLinear { static void call(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { if (stream) call_resize_linear_glob(src, dst, fy, fx, stream); else { if (fx > 1 || fy > 1) call_resize_linear_glob(src, dst, fy, fx, 0); else call_resize_linear_tex(src, srcWhole, yoff, xoff, dst, fy, fx); } } }; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; template <> struct ResizeLinearDispatcher : SelectImplForLinear {}; // ResizeCubicDispatcher template struct ResizeCubicDispatcher { static void call(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { call_resize_cubic_glob(src, dst, fy, fx, stream); } }; template struct SelectImplForCubic { static void call(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream) { if (stream) call_resize_cubic_glob(src, dst, fy, fx, stream); else call_resize_cubic_tex(src, srcWhole, yoff, xoff, dst, fy, fx); } }; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; template <> struct ResizeCubicDispatcher : SelectImplForCubic {}; // ResizeAreaDispatcher template struct ResizeAreaDispatcher { static void call(const PtrStepSz& src, const PtrStepSz&, int, int, const PtrStepSz& dst, float fy, float fx, 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 brd(src.rows, src.cols); BorderReader< PtrStep, BrdConstant > brdSrc(src, brd); IntegerAreaFilter< BorderReader< PtrStep, BrdConstant > > filteredSrc(brdSrc, fx, fy); resize_area<<>>(filteredSrc, dst); } else { BrdConstant brd(src.rows, src.cols); BorderReader< PtrStep, BrdConstant > brdSrc(src, brd); AreaFilter< BorderReader< PtrStep, BrdConstant > > filteredSrc(brdSrc, fx, fy); resize_area<<>>(filteredSrc, dst); } cudaSafeCall( cudaGetLastError() ); if (stream == 0) cudaSafeCall( cudaDeviceSynchronize() ); } }; // resize template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream) { typedef void (*func_t)(const PtrStepSz& src, const PtrStepSz& srcWhole, int yoff, int xoff, const PtrStepSz& dst, float fy, float fx, cudaStream_t stream); static const func_t funcs[4] = { ResizeNearestDispatcher::call, ResizeLinearDispatcher::call, ResizeCubicDispatcher::call, ResizeAreaDispatcher::call }; // change to linear if area interpolation upscaling if (interpolation == 3 && (fx <= 1.f || fy <= 1.f)) interpolation = 1; funcs[interpolation](static_cast< PtrStepSz >(src), static_cast< PtrStepSz >(srcWhole), yoff, xoff, static_cast< PtrStepSz >(dst), fy, fx, stream); } template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); template void resize(const PtrStepSzb& src, const PtrStepSzb& srcWhole, int yoff, int xoff, const PtrStepSzb& dst, float fy, float fx, int interpolation, cudaStream_t stream); }}} #endif /* CUDA_DISABLER */