Skip to content

O
opencv
Commit b2079d6d authored by Vladislav Vinogradov's avatar Vladislav Vinogradov
Browse files
Options

implemented gpu::resize for all types

parent acac27d8
Hide whitespace changes
Inline Side-by-side
Showing with 401 additions and 37 deletions
modules/gpu/include/opencv2/gpu/gpu.hpp
View file @ b2079d6d
......@@ -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
......
modules/gpu/perf/perf_imgproc.cpp
View file @ b2079d6d
......@@ -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()
{
......
modules/gpu/src/cuda/imgproc.cu
View file @ b2079d6d
......@@ -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;
......
modules/gpu/src/imgproc.cpp
View file @ b2079d6d
......@@ -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() );
}
////////////////////////////////////////////////////////////////////////
......
modules/gpu/test/test_imgproc.cpp
View file @ b2079d6d
......@@ -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
......
samples/gpu/performance/tests.cpp
View file @ b2079d6d
......@@ -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;
......@@ -604,6 +676,60 @@ TEST(resize)
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, 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;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment