implemented gpu::addWeighted
This commit is contained in:
Vladislav Vinogradov
parent
64119dd924
commit
0aaaad1ea8
@ -593,6 +593,10 @@ namespace cv
|
||||
//! computes per-element maximum of array and scalar (dst = max(src1, src2))
|
||||
CV_EXPORTS void max(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream = Stream::Null());
|
||||
|
||||
//! computes the weighted sum of two arrays
|
||||
CV_EXPORTS void addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst,
|
||||
int dtype = -1, Stream& stream = Stream::Null());
|
||||
|
||||
|
||||
////////////////////////////// Image processing //////////////////////////////
|
||||
|
||||
|
||||
@ -685,3 +685,34 @@ PERF_TEST_P(DevInfo_Size_MatType, countNonZero, testing::Combine(testing::Values
|
||||
|
||||
SANITY_CHECK(dst);
|
||||
}
|
||||
|
||||
PERF_TEST_P(DevInfo_Size_MatType, addWeighted, testing::Combine(testing::ValuesIn(devices()),
|
||||
testing::Values(GPU_TYPICAL_MAT_SIZES),
|
||||
testing::Values(CV_8UC1, CV_16UC1, CV_32FC1)))
|
||||
{
|
||||
DeviceInfo devInfo = std::tr1::get<0>(GetParam());
|
||||
Size size = std::tr1::get<1>(GetParam());
|
||||
int type = std::tr1::get<2>(GetParam());
|
||||
|
||||
setDevice(devInfo.deviceID());
|
||||
|
||||
Mat src1_host(size, type);
|
||||
Mat src2_host(size, type);
|
||||
|
||||
declare.in(src1_host, src2_host, WARMUP_RNG);
|
||||
|
||||
GpuMat src1(src1_host);
|
||||
GpuMat src2(src2_host);
|
||||
GpuMat dst(size, type);
|
||||
|
||||
declare.time(0.5).iterations(100);
|
||||
|
||||
SIMPLE_TEST_CYCLE()
|
||||
{
|
||||
addWeighted(src1, 0.5, src2, 0.5, 0.0, dst);
|
||||
}
|
||||
|
||||
Mat dst_host = dst;
|
||||
|
||||
SANITY_CHECK(dst_host);
|
||||
}
|
||||
|
||||
@ -757,4 +757,326 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
|
||||
template void multiplyScalar_gpu<uchar, uchar>(const DevMem2D& src, float scale, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
|
||||
//////////////////////////////////////////////////////////////////////////
|
||||
// addWeighted
|
||||
|
||||
template <typename T1, typename T2, typename D> struct AddWeighted : binary_function<T1, T2, D>
|
||||
{
|
||||
__host__ __device__ __forceinline__ AddWeighted(double alpha_, double beta_, double gamma_) : alpha(alpha_), beta(beta_), gamma(gamma_) {}
|
||||
|
||||
__device__ __forceinline__ D operator ()(typename TypeTraits<T1>::ParameterType a, typename TypeTraits<T2>::ParameterType b) const
|
||||
{
|
||||
return saturate_cast<D>(alpha * a + beta * b + gamma);
|
||||
}
|
||||
|
||||
const double alpha;
|
||||
const double beta;
|
||||
const double gamma;
|
||||
};
|
||||
|
||||
template <> struct TransformFunctorTraits< AddWeighted<ushort, ushort, ushort> > : DefaultTransformFunctorTraits< AddWeighted<ushort, ushort, ushort> >
|
||||
{
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<ushort, ushort, short> > : DefaultTransformFunctorTraits< AddWeighted<ushort, ushort, short> >
|
||||
{
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<ushort, short, ushort> > : DefaultTransformFunctorTraits< AddWeighted<ushort, short, ushort> >
|
||||
{
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<ushort, short, short> > : DefaultTransformFunctorTraits< AddWeighted<ushort, short, short> >
|
||||
{
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<short, short, ushort> > : DefaultTransformFunctorTraits< AddWeighted<short, short, ushort> >
|
||||
{
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<short, short, short> > : DefaultTransformFunctorTraits< AddWeighted<short, short, short> >
|
||||
{
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
|
||||
template <> struct TransformFunctorTraits< AddWeighted<int, int, int> > : DefaultTransformFunctorTraits< AddWeighted<int, int, int> >
|
||||
{
|
||||
enum { smart_block_dim_y = 8 };
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<int, int, float> > : DefaultTransformFunctorTraits< AddWeighted<int, int, float> >
|
||||
{
|
||||
enum { smart_block_dim_y = 8 };
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<int, float, int> > : DefaultTransformFunctorTraits< AddWeighted<int, float, int> >
|
||||
{
|
||||
enum { smart_block_dim_y = 8 };
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<int, float, float> > : DefaultTransformFunctorTraits< AddWeighted<int, float, float> >
|
||||
{
|
||||
enum { smart_block_dim_y = 8 };
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<float, float, int> > : DefaultTransformFunctorTraits< AddWeighted<float, float, float> >
|
||||
{
|
||||
enum { smart_block_dim_y = 8 };
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
template <> struct TransformFunctorTraits< AddWeighted<float, float, float> > : DefaultTransformFunctorTraits< AddWeighted<float, float, float> >
|
||||
{
|
||||
enum { smart_block_dim_y = 8 };
|
||||
enum { smart_shift = 4 };
|
||||
};
|
||||
|
||||
template <typename T1, typename T2, typename D>
|
||||
void addWeighted_gpu(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream)
|
||||
{
|
||||
cudaSafeCall( cudaSetDoubleForDevice(&alpha) );
|
||||
cudaSafeCall( cudaSetDoubleForDevice(&beta) );
|
||||
cudaSafeCall( cudaSetDoubleForDevice(&gamma) );
|
||||
|
||||
AddWeighted<T1, T2, D> op(alpha, beta, gamma);
|
||||
|
||||
transform(static_cast< DevMem2D_<T1> >(src1), static_cast< DevMem2D_<T2> >(src2), static_cast< DevMem2D_<D> >(dst), op, stream);
|
||||
}
|
||||
|
||||
template void addWeighted_gpu<uchar, uchar, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, uchar, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, uchar, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, uchar, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, uchar, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, uchar, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, uchar, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<uchar, schar, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, schar, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, schar, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, schar, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, schar, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, schar, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, schar, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<uchar, ushort, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, ushort, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, ushort, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, ushort, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, ushort, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, ushort, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, ushort, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<uchar, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<uchar, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<uchar, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<uchar, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<uchar, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
|
||||
|
||||
template void addWeighted_gpu<schar, schar, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, schar, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, schar, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, schar, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, schar, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, schar, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, schar, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<schar, ushort, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, ushort, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, ushort, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, ushort, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, ushort, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, ushort, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, ushort, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<schar, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<schar, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<schar, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<schar, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<schar, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
|
||||
|
||||
template void addWeighted_gpu<ushort, ushort, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, ushort, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, ushort, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, ushort, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, ushort, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, ushort, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, ushort, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<ushort, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<ushort, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<ushort, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<ushort, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<ushort, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
|
||||
|
||||
template void addWeighted_gpu<short, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<short, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<short, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<short, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<short, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
|
||||
|
||||
template void addWeighted_gpu<int, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<int, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<int, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<int, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
|
||||
|
||||
template void addWeighted_gpu<float, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
template void addWeighted_gpu<float, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<float, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
|
||||
|
||||
template void addWeighted_gpu<double, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<double, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<double, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<double, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<double, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<double, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
template void addWeighted_gpu<double, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
}}}
|
||||
|
||||
@ -67,7 +67,8 @@ void cv::gpu::min(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }
|
||||
void cv::gpu::max(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
|
||||
void cv::gpu::max(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }
|
||||
double cv::gpu::threshold(const GpuMat&, GpuMat&, double, double, int, Stream&) {throw_nogpu(); return 0.0;}
|
||||
void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }
|
||||
void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }
|
||||
void cv::gpu::addWeighted(const GpuMat&, double, const GpuMat&, double, double, GpuMat&, int, Stream&) { throw_nogpu(); }
|
||||
|
||||
#else
|
||||
|
||||
@ -248,7 +249,7 @@ void cv::gpu::multiply(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream&
|
||||
{
|
||||
dst.create(src.size(), src.type());
|
||||
|
||||
device::multiplyScalar_gpu<uchar, uchar>(src.reshape(1), (float)(sc[0]), dst, StreamAccessor::getStream(stream));
|
||||
device::multiplyScalar_gpu<unsigned char, unsigned char>(src.reshape(1), (float)(sc[0]), dst, StreamAccessor::getStream(stream));
|
||||
}
|
||||
else
|
||||
{
|
||||
@ -733,7 +734,7 @@ void cv::gpu::min(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& s
|
||||
typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream);
|
||||
static const func_t funcs[] =
|
||||
{
|
||||
min_caller<uchar>, min_caller<schar>, min_caller<ushort>, min_caller<short>, min_caller<int>,
|
||||
min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>,
|
||||
min_caller<float>, min_caller<double>
|
||||
};
|
||||
funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
|
||||
@ -746,7 +747,7 @@ void cv::gpu::min(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream)
|
||||
typedef void (*func_t)(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream);
|
||||
static const func_t funcs[] =
|
||||
{
|
||||
min_caller<uchar>, min_caller<schar>, min_caller<ushort>, min_caller<short>, min_caller<int>,
|
||||
min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>,
|
||||
min_caller<float>, min_caller<double>
|
||||
};
|
||||
funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
|
||||
@ -761,7 +762,7 @@ void cv::gpu::max(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& s
|
||||
typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream);
|
||||
static const func_t funcs[] =
|
||||
{
|
||||
max_caller<uchar>, max_caller<schar>, max_caller<ushort>, max_caller<short>, max_caller<int>,
|
||||
max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>,
|
||||
max_caller<float>, max_caller<double>
|
||||
};
|
||||
funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
|
||||
@ -775,7 +776,7 @@ void cv::gpu::max(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream)
|
||||
typedef void (*func_t)(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream);
|
||||
static const func_t funcs[] =
|
||||
{
|
||||
max_caller<uchar>, max_caller<schar>, max_caller<ushort>, max_caller<short>, max_caller<int>,
|
||||
max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>,
|
||||
max_caller<float>, max_caller<double>
|
||||
};
|
||||
funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
|
||||
@ -875,4 +876,497 @@ void cv::gpu::pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream)
|
||||
callers[src.depth()](src.reshape(1), (float)power, dst.reshape(1), StreamAccessor::getStream(stream));
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
// addWeighted
|
||||
|
||||
namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
template <typename T1, typename T2, typename D>
|
||||
void addWeighted_gpu(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
}}}
|
||||
|
||||
void cv::gpu::addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, int dtype, Stream& stream)
|
||||
{
|
||||
CV_Assert(src1.size() == src2.size());
|
||||
CV_Assert(src1.type() == src2.type() || (dtype >= 0 && src1.channels() == src2.channels()));
|
||||
|
||||
dtype = dtype >= 0 ? CV_MAKETYPE(dtype, src1.channels()) : src1.type();
|
||||
|
||||
dst.create(src1.size(), dtype);
|
||||
|
||||
const GpuMat* psrc1 = &src1;
|
||||
const GpuMat* psrc2 = &src2;
|
||||
|
||||
if (src1.depth() > src2.depth())
|
||||
{
|
||||
std::swap(psrc1, psrc2);
|
||||
std::swap(alpha, beta);
|
||||
}
|
||||
|
||||
typedef void (*caller_t)(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
|
||||
|
||||
using namespace cv::gpu::device;
|
||||
|
||||
static const caller_t callers[7][7][7] =
|
||||
{
|
||||
{
|
||||
{
|
||||
addWeighted_gpu<unsigned char, unsigned char, unsigned char >,
|
||||
addWeighted_gpu<unsigned char, unsigned char, signed char >,
|
||||
addWeighted_gpu<unsigned char, unsigned char, unsigned short>,
|
||||
addWeighted_gpu<unsigned char, unsigned char, short >,
|
||||
addWeighted_gpu<unsigned char, unsigned char, int >,
|
||||
addWeighted_gpu<unsigned char, unsigned char, float >,
|
||||
addWeighted_gpu<unsigned char, unsigned char, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned char, signed char, unsigned char >,
|
||||
addWeighted_gpu<unsigned char, signed char, signed char >,
|
||||
addWeighted_gpu<unsigned char, signed char, unsigned short>,
|
||||
addWeighted_gpu<unsigned char, signed char, short >,
|
||||
addWeighted_gpu<unsigned char, signed char, int >,
|
||||
addWeighted_gpu<unsigned char, signed char, float >,
|
||||
addWeighted_gpu<unsigned char, signed char, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned char, unsigned short, unsigned char >,
|
||||
addWeighted_gpu<unsigned char, unsigned short, signed char >,
|
||||
addWeighted_gpu<unsigned char, unsigned short, unsigned short>,
|
||||
addWeighted_gpu<unsigned char, unsigned short, short >,
|
||||
addWeighted_gpu<unsigned char, unsigned short, int >,
|
||||
addWeighted_gpu<unsigned char, unsigned short, float >,
|
||||
addWeighted_gpu<unsigned char, unsigned short, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned char, short, unsigned char >,
|
||||
addWeighted_gpu<unsigned char, short, signed char >,
|
||||
addWeighted_gpu<unsigned char, short, unsigned short>,
|
||||
addWeighted_gpu<unsigned char, short, short >,
|
||||
addWeighted_gpu<unsigned char, short, int >,
|
||||
addWeighted_gpu<unsigned char, short, float >,
|
||||
addWeighted_gpu<unsigned char, short, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned char, int, unsigned char >,
|
||||
addWeighted_gpu<unsigned char, int, signed char >,
|
||||
addWeighted_gpu<unsigned char, int, unsigned short>,
|
||||
addWeighted_gpu<unsigned char, int, short >,
|
||||
addWeighted_gpu<unsigned char, int, int >,
|
||||
addWeighted_gpu<unsigned char, int, float >,
|
||||
addWeighted_gpu<unsigned char, int, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned char, float, unsigned char >,
|
||||
addWeighted_gpu<unsigned char, float, signed char >,
|
||||
addWeighted_gpu<unsigned char, float, unsigned short>,
|
||||
addWeighted_gpu<unsigned char, float, short >,
|
||||
addWeighted_gpu<unsigned char, float, int >,
|
||||
addWeighted_gpu<unsigned char, float, float >,
|
||||
addWeighted_gpu<unsigned char, float, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned char, double, unsigned char >,
|
||||
addWeighted_gpu<unsigned char, double, signed char >,
|
||||
addWeighted_gpu<unsigned char, double, unsigned short>,
|
||||
addWeighted_gpu<unsigned char, double, short >,
|
||||
addWeighted_gpu<unsigned char, double, int >,
|
||||
addWeighted_gpu<unsigned char, double, float >,
|
||||
addWeighted_gpu<unsigned char, double, double>
|
||||
}
|
||||
},
|
||||
{
|
||||
{
|
||||
0/*addWeighted_gpu<signed char, unsigned char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<signed char, unsigned char, signed char >*/,
|
||||
0/*addWeighted_gpu<signed char, unsigned char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<signed char, unsigned char, short >*/,
|
||||
0/*addWeighted_gpu<signed char, unsigned char, int >*/,
|
||||
0/*addWeighted_gpu<signed char, unsigned char, float >*/,
|
||||
0/*addWeighted_gpu<signed char, unsigned char, double>*/
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<signed char, signed char, unsigned char >,
|
||||
addWeighted_gpu<signed char, signed char, signed char >,
|
||||
addWeighted_gpu<signed char, signed char, unsigned short>,
|
||||
addWeighted_gpu<signed char, signed char, short >,
|
||||
addWeighted_gpu<signed char, signed char, int >,
|
||||
addWeighted_gpu<signed char, signed char, float >,
|
||||
addWeighted_gpu<signed char, signed char, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<signed char, unsigned short, unsigned char >,
|
||||
addWeighted_gpu<signed char, unsigned short, signed char >,
|
||||
addWeighted_gpu<signed char, unsigned short, unsigned short>,
|
||||
addWeighted_gpu<signed char, unsigned short, short >,
|
||||
addWeighted_gpu<signed char, unsigned short, int >,
|
||||
addWeighted_gpu<signed char, unsigned short, float >,
|
||||
addWeighted_gpu<signed char, unsigned short, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<signed char, short, unsigned char >,
|
||||
addWeighted_gpu<signed char, short, signed char >,
|
||||
addWeighted_gpu<signed char, short, unsigned short>,
|
||||
addWeighted_gpu<signed char, short, short >,
|
||||
addWeighted_gpu<signed char, short, int >,
|
||||
addWeighted_gpu<signed char, short, float >,
|
||||
addWeighted_gpu<signed char, short, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<signed char, int, unsigned char >,
|
||||
addWeighted_gpu<signed char, int, signed char >,
|
||||
addWeighted_gpu<signed char, int, unsigned short>,
|
||||
addWeighted_gpu<signed char, int, short >,
|
||||
addWeighted_gpu<signed char, int, int >,
|
||||
addWeighted_gpu<signed char, int, float >,
|
||||
addWeighted_gpu<signed char, int, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<signed char, float, unsigned char >,
|
||||
addWeighted_gpu<signed char, float, signed char >,
|
||||
addWeighted_gpu<signed char, float, unsigned short>,
|
||||
addWeighted_gpu<signed char, float, short >,
|
||||
addWeighted_gpu<signed char, float, int >,
|
||||
addWeighted_gpu<signed char, float, float >,
|
||||
addWeighted_gpu<signed char, float, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<signed char, double, unsigned char >,
|
||||
addWeighted_gpu<signed char, double, signed char >,
|
||||
addWeighted_gpu<signed char, double, unsigned short>,
|
||||
addWeighted_gpu<signed char, double, short >,
|
||||
addWeighted_gpu<signed char, double, int >,
|
||||
addWeighted_gpu<signed char, double, float >,
|
||||
addWeighted_gpu<signed char, double, double>
|
||||
}
|
||||
},
|
||||
{
|
||||
{
|
||||
0/*addWeighted_gpu<unsigned short, unsigned char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<unsigned short, unsigned char, signed char >*/,
|
||||
0/*addWeighted_gpu<unsigned short, unsigned char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<unsigned short, unsigned char, short >*/,
|
||||
0/*addWeighted_gpu<unsigned short, unsigned char, int >*/,
|
||||
0/*addWeighted_gpu<unsigned short, unsigned char, float >*/,
|
||||
0/*addWeighted_gpu<unsigned short, unsigned char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<unsigned short, signed char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<unsigned short, signed char, signed char >*/,
|
||||
0/*addWeighted_gpu<unsigned short, signed char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<unsigned short, signed char, short >*/,
|
||||
0/*addWeighted_gpu<unsigned short, signed char, int >*/,
|
||||
0/*addWeighted_gpu<unsigned short, signed char, float >*/,
|
||||
0/*addWeighted_gpu<unsigned short, signed char, double>*/
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned short, unsigned short, unsigned char >,
|
||||
addWeighted_gpu<unsigned short, unsigned short, signed char >,
|
||||
addWeighted_gpu<unsigned short, unsigned short, unsigned short>,
|
||||
addWeighted_gpu<unsigned short, unsigned short, short >,
|
||||
addWeighted_gpu<unsigned short, unsigned short, int >,
|
||||
addWeighted_gpu<unsigned short, unsigned short, float >,
|
||||
addWeighted_gpu<unsigned short, unsigned short, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned short, short, unsigned char >,
|
||||
addWeighted_gpu<unsigned short, short, signed char >,
|
||||
addWeighted_gpu<unsigned short, short, unsigned short>,
|
||||
addWeighted_gpu<unsigned short, short, short >,
|
||||
addWeighted_gpu<unsigned short, short, int >,
|
||||
addWeighted_gpu<unsigned short, short, float >,
|
||||
addWeighted_gpu<unsigned short, short, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned short, int, unsigned char >,
|
||||
addWeighted_gpu<unsigned short, int, signed char >,
|
||||
addWeighted_gpu<unsigned short, int, unsigned short>,
|
||||
addWeighted_gpu<unsigned short, int, short >,
|
||||
addWeighted_gpu<unsigned short, int, int >,
|
||||
addWeighted_gpu<unsigned short, int, float >,
|
||||
addWeighted_gpu<unsigned short, int, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned short, float, unsigned char >,
|
||||
addWeighted_gpu<unsigned short, float, signed char >,
|
||||
addWeighted_gpu<unsigned short, float, unsigned short>,
|
||||
addWeighted_gpu<unsigned short, float, short >,
|
||||
addWeighted_gpu<unsigned short, float, int >,
|
||||
addWeighted_gpu<unsigned short, float, float >,
|
||||
addWeighted_gpu<unsigned short, float, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<unsigned short, double, unsigned char >,
|
||||
addWeighted_gpu<unsigned short, double, signed char >,
|
||||
addWeighted_gpu<unsigned short, double, unsigned short>,
|
||||
addWeighted_gpu<unsigned short, double, short >,
|
||||
addWeighted_gpu<unsigned short, double, int >,
|
||||
addWeighted_gpu<unsigned short, double, float >,
|
||||
addWeighted_gpu<unsigned short, double, double>
|
||||
}
|
||||
},
|
||||
{
|
||||
{
|
||||
0/*addWeighted_gpu<short, unsigned char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<short, unsigned char, signed char >*/,
|
||||
0/*addWeighted_gpu<short, unsigned char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<short, unsigned char, short >*/,
|
||||
0/*addWeighted_gpu<short, unsigned char, int >*/,
|
||||
0/*addWeighted_gpu<short, unsigned char, float >*/,
|
||||
0/*addWeighted_gpu<short, unsigned char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<short, signed char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<short, signed char, signed char >*/,
|
||||
0/*addWeighted_gpu<short, signed char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<short, signed char, short >*/,
|
||||
0/*addWeighted_gpu<short, signed char, int >*/,
|
||||
0/*addWeighted_gpu<short, signed char, float >*/,
|
||||
0/*addWeighted_gpu<short, signed char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<short, unsigned short, unsigned char >*/,
|
||||
0/*addWeighted_gpu<short, unsigned short, signed char >*/,
|
||||
0/*addWeighted_gpu<short, unsigned short, unsigned short>*/,
|
||||
0/*addWeighted_gpu<short, unsigned short, short >*/,
|
||||
0/*addWeighted_gpu<short, unsigned short, int >*/,
|
||||
0/*addWeighted_gpu<short, unsigned short, float >*/,
|
||||
0/*addWeighted_gpu<short, unsigned short, double>*/
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<short, short, unsigned char >,
|
||||
addWeighted_gpu<short, short, signed char >,
|
||||
addWeighted_gpu<short, short, unsigned short>,
|
||||
addWeighted_gpu<short, short, short >,
|
||||
addWeighted_gpu<short, short, int >,
|
||||
addWeighted_gpu<short, short, float >,
|
||||
addWeighted_gpu<short, short, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<short, int, unsigned char >,
|
||||
addWeighted_gpu<short, int, signed char >,
|
||||
addWeighted_gpu<short, int, unsigned short>,
|
||||
addWeighted_gpu<short, int, short >,
|
||||
addWeighted_gpu<short, int, int >,
|
||||
addWeighted_gpu<short, int, float >,
|
||||
addWeighted_gpu<short, int, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<short, float, unsigned char >,
|
||||
addWeighted_gpu<short, float, signed char >,
|
||||
addWeighted_gpu<short, float, unsigned short>,
|
||||
addWeighted_gpu<short, float, short >,
|
||||
addWeighted_gpu<short, float, int >,
|
||||
addWeighted_gpu<short, float, float >,
|
||||
addWeighted_gpu<short, float, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<short, double, unsigned char >,
|
||||
addWeighted_gpu<short, double, signed char >,
|
||||
addWeighted_gpu<short, double, unsigned short>,
|
||||
addWeighted_gpu<short, double, short >,
|
||||
addWeighted_gpu<short, double, int >,
|
||||
addWeighted_gpu<short, double, float >,
|
||||
addWeighted_gpu<short, double, double>
|
||||
}
|
||||
},
|
||||
{
|
||||
{
|
||||
0/*addWeighted_gpu<int, unsigned char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<int, unsigned char, signed char >*/,
|
||||
0/*addWeighted_gpu<int, unsigned char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<int, unsigned char, short >*/,
|
||||
0/*addWeighted_gpu<int, unsigned char, int >*/,
|
||||
0/*addWeighted_gpu<int, unsigned char, float >*/,
|
||||
0/*addWeighted_gpu<int, unsigned char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<int, signed char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<int, signed char, signed char >*/,
|
||||
0/*addWeighted_gpu<int, signed char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<int, signed char, short >*/,
|
||||
0/*addWeighted_gpu<int, signed char, int >*/,
|
||||
0/*addWeighted_gpu<int, signed char, float >*/,
|
||||
0/*addWeighted_gpu<int, signed char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<int, unsigned short, unsigned char >*/,
|
||||
0/*addWeighted_gpu<int, unsigned short, signed char >*/,
|
||||
0/*addWeighted_gpu<int, unsigned short, unsigned short>*/,
|
||||
0/*addWeighted_gpu<int, unsigned short, short >*/,
|
||||
0/*addWeighted_gpu<int, unsigned short, int >*/,
|
||||
0/*addWeighted_gpu<int, unsigned short, float >*/,
|
||||
0/*addWeighted_gpu<int, unsigned short, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<int, short, unsigned char >*/,
|
||||
0/*addWeighted_gpu<int, short, signed char >*/,
|
||||
0/*addWeighted_gpu<int, short, unsigned short>*/,
|
||||
0/*addWeighted_gpu<int, short, short >*/,
|
||||
0/*addWeighted_gpu<int, short, int >*/,
|
||||
0/*addWeighted_gpu<int, short, float >*/,
|
||||
0/*addWeighted_gpu<int, short, double>*/
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<int, int, unsigned char >,
|
||||
addWeighted_gpu<int, int, signed char >,
|
||||
addWeighted_gpu<int, int, unsigned short>,
|
||||
addWeighted_gpu<int, int, short >,
|
||||
addWeighted_gpu<int, int, int >,
|
||||
addWeighted_gpu<int, int, float >,
|
||||
addWeighted_gpu<int, int, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<int, float, unsigned char >,
|
||||
addWeighted_gpu<int, float, signed char >,
|
||||
addWeighted_gpu<int, float, unsigned short>,
|
||||
addWeighted_gpu<int, float, short >,
|
||||
addWeighted_gpu<int, float, int >,
|
||||
addWeighted_gpu<int, float, float >,
|
||||
addWeighted_gpu<int, float, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<int, double, unsigned char >,
|
||||
addWeighted_gpu<int, double, signed char >,
|
||||
addWeighted_gpu<int, double, unsigned short>,
|
||||
addWeighted_gpu<int, double, short >,
|
||||
addWeighted_gpu<int, double, int >,
|
||||
addWeighted_gpu<int, double, float >,
|
||||
addWeighted_gpu<int, double, double>
|
||||
}
|
||||
},
|
||||
{
|
||||
{
|
||||
0/*addWeighted_gpu<float, unsigned char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<float, unsigned char, signed char >*/,
|
||||
0/*addWeighted_gpu<float, unsigned char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<float, unsigned char, short >*/,
|
||||
0/*addWeighted_gpu<float, unsigned char, int >*/,
|
||||
0/*addWeighted_gpu<float, unsigned char, float >*/,
|
||||
0/*addWeighted_gpu<float, unsigned char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<float, signed char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<float, signed char, signed char >*/,
|
||||
0/*addWeighted_gpu<float, signed char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<float, signed char, short >*/,
|
||||
0/*addWeighted_gpu<float, signed char, int >*/,
|
||||
0/*addWeighted_gpu<float, signed char, float >*/,
|
||||
0/*addWeighted_gpu<float, signed char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<float, unsigned short, unsigned char >*/,
|
||||
0/*addWeighted_gpu<float, unsigned short, signed char >*/,
|
||||
0/*addWeighted_gpu<float, unsigned short, unsigned short>*/,
|
||||
0/*addWeighted_gpu<float, unsigned short, short >*/,
|
||||
0/*addWeighted_gpu<float, unsigned short, int >*/,
|
||||
0/*addWeighted_gpu<float, unsigned short, float >*/,
|
||||
0/*addWeighted_gpu<float, unsigned short, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<float, short, unsigned char >*/,
|
||||
0/*addWeighted_gpu<float, short, signed char >*/,
|
||||
0/*addWeighted_gpu<float, short, unsigned short>*/,
|
||||
0/*addWeighted_gpu<float, short, short >*/,
|
||||
0/*addWeighted_gpu<float, short, int >*/,
|
||||
0/*addWeighted_gpu<float, short, float >*/,
|
||||
0/*addWeighted_gpu<float, short, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<float, int, unsigned char >*/,
|
||||
0/*addWeighted_gpu<float, int, signed char >*/,
|
||||
0/*addWeighted_gpu<float, int, unsigned short>*/,
|
||||
0/*addWeighted_gpu<float, int, short >*/,
|
||||
0/*addWeighted_gpu<float, int, int >*/,
|
||||
0/*addWeighted_gpu<float, int, float >*/,
|
||||
0/*addWeighted_gpu<float, int, double>*/
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<float, float, unsigned char >,
|
||||
addWeighted_gpu<float, float, signed char >,
|
||||
addWeighted_gpu<float, float, unsigned short>,
|
||||
addWeighted_gpu<float, float, short >,
|
||||
addWeighted_gpu<float, float, int >,
|
||||
addWeighted_gpu<float, float, float >,
|
||||
addWeighted_gpu<float, float, double>
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<float, double, unsigned char >,
|
||||
addWeighted_gpu<float, double, signed char >,
|
||||
addWeighted_gpu<float, double, unsigned short>,
|
||||
addWeighted_gpu<float, double, short >,
|
||||
addWeighted_gpu<float, double, int >,
|
||||
addWeighted_gpu<float, double, float >,
|
||||
addWeighted_gpu<float, double, double>
|
||||
}
|
||||
},
|
||||
{
|
||||
{
|
||||
0/*addWeighted_gpu<double, unsigned char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<double, unsigned char, signed char >*/,
|
||||
0/*addWeighted_gpu<double, unsigned char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<double, unsigned char, short >*/,
|
||||
0/*addWeighted_gpu<double, unsigned char, int >*/,
|
||||
0/*addWeighted_gpu<double, unsigned char, float >*/,
|
||||
0/*addWeighted_gpu<double, unsigned char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<double, signed char, unsigned char >*/,
|
||||
0/*addWeighted_gpu<double, signed char, signed char >*/,
|
||||
0/*addWeighted_gpu<double, signed char, unsigned short>*/,
|
||||
0/*addWeighted_gpu<double, signed char, short >*/,
|
||||
0/*addWeighted_gpu<double, signed char, int >*/,
|
||||
0/*addWeighted_gpu<double, signed char, float >*/,
|
||||
0/*addWeighted_gpu<double, signed char, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<double, unsigned short, unsigned char >*/,
|
||||
0/*addWeighted_gpu<double, unsigned short, signed char >*/,
|
||||
0/*addWeighted_gpu<double, unsigned short, unsigned short>*/,
|
||||
0/*addWeighted_gpu<double, unsigned short, short >*/,
|
||||
0/*addWeighted_gpu<double, unsigned short, int >*/,
|
||||
0/*addWeighted_gpu<double, unsigned short, float >*/,
|
||||
0/*addWeighted_gpu<double, unsigned short, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<double, short, unsigned char >*/,
|
||||
0/*addWeighted_gpu<double, short, signed char >*/,
|
||||
0/*addWeighted_gpu<double, short, unsigned short>*/,
|
||||
0/*addWeighted_gpu<double, short, short >*/,
|
||||
0/*addWeighted_gpu<double, short, int >*/,
|
||||
0/*addWeighted_gpu<double, short, float >*/,
|
||||
0/*addWeighted_gpu<double, short, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<double, int, unsigned char >*/,
|
||||
0/*addWeighted_gpu<double, int, signed char >*/,
|
||||
0/*addWeighted_gpu<double, int, unsigned short>*/,
|
||||
0/*addWeighted_gpu<double, int, short >*/,
|
||||
0/*addWeighted_gpu<double, int, int >*/,
|
||||
0/*addWeighted_gpu<double, int, float >*/,
|
||||
0/*addWeighted_gpu<double, int, double>*/
|
||||
},
|
||||
{
|
||||
0/*addWeighted_gpu<double, float, unsigned char >*/,
|
||||
0/*addWeighted_gpu<double, float, signed char >*/,
|
||||
0/*addWeighted_gpu<double, float, unsigned short>*/,
|
||||
0/*addWeighted_gpu<double, float, short >*/,
|
||||
0/*addWeighted_gpu<double, float, int >*/,
|
||||
0/*addWeighted_gpu<double, float, float >*/,
|
||||
0/*addWeighted_gpu<double, float, double>*/
|
||||
},
|
||||
{
|
||||
addWeighted_gpu<double, double, unsigned char >,
|
||||
addWeighted_gpu<double, double, signed char >,
|
||||
addWeighted_gpu<double, double, unsigned short>,
|
||||
addWeighted_gpu<double, double, short >,
|
||||
addWeighted_gpu<double, double, int >,
|
||||
addWeighted_gpu<double, double, float >,
|
||||
addWeighted_gpu<double, double, double>
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
callers[psrc1->depth()][psrc2->depth()][dst.depth()](psrc1->reshape(1), alpha, psrc2->reshape(1), beta, gamma, dst.reshape(1), StreamAccessor::getStream(stream));
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
@ -1135,7 +1135,7 @@ TEST_P(MinMax, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
double minVal, maxVal;
|
||||
@ -1216,7 +1216,7 @@ TEST_P(MinMaxLoc, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
double minVal, maxVal;
|
||||
@ -1281,7 +1281,7 @@ TEST_P(CountNonZero, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
int n;
|
||||
@ -1333,7 +1333,7 @@ TEST_P(Sum, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
cv::Scalar sum;
|
||||
@ -1385,7 +1385,7 @@ TEST_P(AbsSum, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
cv::Scalar sum;
|
||||
@ -1439,7 +1439,7 @@ TEST_P(SqrSum, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
cv::Scalar sum;
|
||||
@ -1500,7 +1500,7 @@ TEST_P(BitwiseNot, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
cv::Mat dst;
|
||||
@ -1564,7 +1564,7 @@ TEST_P(BitwiseOr, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
cv::Mat dst;
|
||||
@ -1628,7 +1628,7 @@ TEST_P(BitwiseAnd, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
cv::Mat dst;
|
||||
@ -1692,7 +1692,7 @@ TEST_P(BitwiseXor, Accuracy)
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type)
|
||||
PRINT_TYPE(type);
|
||||
PRINT_PARAM(size);
|
||||
|
||||
cv::Mat dst;
|
||||
@ -1712,4 +1712,80 @@ INSTANTIATE_TEST_CASE_P(Arithm, BitwiseXor, testing::Combine(
|
||||
testing::ValuesIn(devices()),
|
||||
testing::ValuesIn(all_types())));
|
||||
|
||||
//////////////////////////////////////////////////////////////////////////////
|
||||
// addWeighted
|
||||
|
||||
struct AddWeighted : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int, int, int> >
|
||||
{
|
||||
cv::gpu::DeviceInfo devInfo;
|
||||
int type1;
|
||||
int type2;
|
||||
int dtype;
|
||||
|
||||
cv::Size size;
|
||||
cv::Mat src1;
|
||||
cv::Mat src2;
|
||||
double alpha;
|
||||
double beta;
|
||||
double gamma;
|
||||
|
||||
cv::Mat dst_gold;
|
||||
|
||||
virtual void SetUp()
|
||||
{
|
||||
devInfo = std::tr1::get<0>(GetParam());
|
||||
type1 = std::tr1::get<1>(GetParam());
|
||||
type2 = std::tr1::get<2>(GetParam());
|
||||
dtype = std::tr1::get<3>(GetParam());
|
||||
|
||||
cv::gpu::setDevice(devInfo.deviceID());
|
||||
|
||||
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
|
||||
|
||||
size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200));
|
||||
|
||||
src1 = cvtest::randomMat(rng, size, type1, 0.0, 255.0, false);
|
||||
src2 = cvtest::randomMat(rng, size, type2, 0.0, 255.0, false);
|
||||
|
||||
alpha = rng.uniform(-10.0, 10.0);
|
||||
beta = rng.uniform(-10.0, 10.0);
|
||||
gamma = rng.uniform(-10.0, 10.0);
|
||||
|
||||
cv::addWeighted(src1, alpha, src2, beta, gamma, dst_gold, dtype);
|
||||
}
|
||||
};
|
||||
|
||||
TEST_P(AddWeighted, Accuracy)
|
||||
{
|
||||
if ((src1.depth() == CV_64F || src2.depth() == CV_64F || dst_gold.depth() == CV_64F) && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
|
||||
return;
|
||||
|
||||
PRINT_PARAM(devInfo);
|
||||
PRINT_TYPE(type1);
|
||||
PRINT_TYPE(type2);
|
||||
PRINT_TYPE(dtype);
|
||||
PRINT_PARAM(size);
|
||||
PRINT_PARAM(alpha);
|
||||
PRINT_PARAM(beta);
|
||||
PRINT_PARAM(gamma);
|
||||
|
||||
cv::Mat dst;
|
||||
|
||||
ASSERT_NO_THROW(
|
||||
cv::gpu::GpuMat dev_dst;
|
||||
|
||||
cv::gpu::addWeighted(cv::gpu::GpuMat(src1), alpha, cv::gpu::GpuMat(src2), beta, gamma, dev_dst, dtype);
|
||||
|
||||
dev_dst.download(dst);
|
||||
);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, dtype < CV_32F ? 1.0 : 1e-12);
|
||||
}
|
||||
|
||||
INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, testing::Combine(
|
||||
testing::ValuesIn(devices()),
|
||||
testing::ValuesIn(types(CV_8U, CV_64F, 1, 1)),
|
||||
testing::ValuesIn(types(CV_8U, CV_64F, 1, 1)),
|
||||
testing::ValuesIn(types(CV_8U, CV_64F, 1, 1))));
|
||||
|
||||
#endif // HAVE_CUDA
|
||||
|
||||
Loading…
Reference in New Issue
Block a user