mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-24 19:20:28 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Use in-place npp function for inplace arguments

Browse Source
This commit is contained in:
Namgoo Lee 2020-07-21 10:27:43 +09:00
parent 284d26da05
commit 9411cd6c07
2 changed files with 56 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

44
modules/cudaarithm/src/core.cpp
View File

@ -102,6 +102,34 @@ namespace
cudaSafeCall( cudaDeviceSynchronize() );
}
};
template <int DEPTH> struct NppMirrorIFunc
{
typedef typename NppTypeTraits<DEPTH>::npp_t npp_t;
typedef NppStatus (*func_t)(npp_t* pSrcDst, int nSrcDstStep, NppiSize oROI, NppiAxis flip);
};
template <int DEPTH, typename NppMirrorIFunc<DEPTH>::func_t func> struct NppMirrorI
{
typedef typename NppMirrorIFunc<DEPTH>::npp_t npp_t;
static void call(GpuMat& srcDst, int flipCode, cudaStream_t stream)
{
NppStreamHandler h(stream);
NppiSize sz;
sz.width = srcDst.cols;
sz.height = srcDst.rows;
nppSafeCall( func(srcDst.ptr<npp_t>(), static_cast<int>(srcDst.step),
sz,
(flipCode == 0 ? NPP_HORIZONTAL_AXIS : (flipCode > 0 ? NPP_VERTICAL_AXIS : NPP_BOTH_AXIS))) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
}
void cv::cuda::flip(InputArray _src, OutputArray _dst, int flipCode, Stream& stream)
@ -117,6 +145,17 @@ void cv::cuda::flip(InputArray _src, OutputArray _dst, int flipCode, Stream& str
{NppMirror<CV_32F, nppiMirror_32f_C1R>::call, 0, NppMirror<CV_32F, nppiMirror_32f_C3R>::call, NppMirror<CV_32F, nppiMirror_32f_C4R>::call}
};
typedef void (*ifunc_t)(GpuMat& srcDst, int flipCode, cudaStream_t stream);
static const ifunc_t ifuncs[6][4] =
{
{NppMirrorI<CV_8U, nppiMirror_8u_C1IR>::call, 0, NppMirrorI<CV_8U, nppiMirror_8u_C3IR>::call, NppMirrorI<CV_8U, nppiMirror_8u_C4IR>::call},
{0,0,0,0},
{NppMirrorI<CV_16U, nppiMirror_16u_C1IR>::call, 0, NppMirrorI<CV_16U, nppiMirror_16u_C3IR>::call, NppMirrorI<CV_16U, nppiMirror_16u_C4IR>::call},
{0,0,0,0},
{NppMirrorI<CV_32S, nppiMirror_32s_C1IR>::call, 0, NppMirrorI<CV_32S, nppiMirror_32s_C3IR>::call, NppMirrorI<CV_32S, nppiMirror_32s_C4IR>::call},
{NppMirrorI<CV_32F, nppiMirror_32f_C1IR>::call, 0, NppMirrorI<CV_32F, nppiMirror_32f_C3IR>::call, NppMirrorI<CV_32F, nppiMirror_32f_C4IR>::call}
};
GpuMat src = getInputMat(_src, stream);
CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F);
@ -125,7 +164,10 @@ void cv::cuda::flip(InputArray _src, OutputArray _dst, int flipCode, Stream& str
_dst.create(src.size(), src.type());
GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
funcs[src.depth()][src.channels() - 1](src, dst, flipCode, StreamAccessor::getStream(stream));
if (src.refcount != dst.refcount)
funcs[src.depth()][src.channels() - 1](src, dst, flipCode, StreamAccessor::getStream(stream));
else // in-place
ifuncs[src.depth()][src.channels() - 1](src, flipCode, StreamAccessor::getStream(stream));
syncOutput(dst, _dst, stream);
}

13
modules/cudaarithm/test/test_core.cpp
View File

@ -279,6 +279,19 @@ CUDA_TEST_P(Flip, Accuracy)
EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
}
CUDA_TEST_P(Flip, AccuracyInplace)
{
cv::Mat src = randomMat(size, type);
cv::cuda::GpuMat srcDst = loadMat(src, useRoi);
cv::cuda::flip(srcDst, srcDst, flip_code);
cv::Mat dst_gold;
cv::flip(src, dst_gold, flip_code);
EXPECT_MAT_NEAR(dst_gold, srcDst, 0.0);
}
INSTANTIATE_TEST_CASE_P(CUDA_Arithm, Flip, testing::Combine(
ALL_DEVICES,
DIFFERENT_SIZES,