Skip to content

O
opencv_contrib
Commit a345402d authored by Alexander Alekhin's avatar Alexander Alekhin
Browse files
Options

Merge moved code from opencv

parents 2e3cbde7 c7c70361
Hide whitespace changes
Inline Side-by-side
Showing with 98 additions and 66 deletions
modules/cudaarithm/include/opencv2/cudaarithm.hpp
View file @ a345402d
......@@ -208,7 +208,7 @@ CV_EXPORTS_W void pow(InputArray src, double power, OutputArray dst, Stream& str
@param src1 First source matrix or scalar.
@param src2 Second source matrix or scalar.
@param dst Destination matrix that has the same size and type as the input array(s).
@param dst Destination matrix that has the same size as the input array(s) and type CV_8U.
@param cmpop Flag specifying the relation between the elements to be checked:
- **CMP_EQ:** a(.) == b(.)
- **CMP_GT:** a(.) \> b(.)
......@@ -415,10 +415,10 @@ CV_EXPORTS_W void cartToPolar(InputArray x, InputArray y, OutputArray magnitude,
/** @brief Converts polar coordinates into Cartesian.
@param magnitude Source matrix containing magnitudes ( CV_32FC1 ).
@param angle Source matrix containing angles ( CV_32FC1 ).
@param x Destination matrix of real components ( CV_32FC1 ).
@param y Destination matrix of imaginary components ( CV_32FC1 ).
@param magnitude Source matrix containing magnitudes ( CV_32FC1 or CV_64FC1 ).
@param angle Source matrix containing angles ( same type as magnitude ).
@param x Destination matrix of real components ( same type as magnitude ).
@param y Destination matrix of imaginary components ( same type as magnitude ).
@param angleInDegrees Flag that indicates angles in degrees.
@param stream Stream for the asynchronous version.
*/
......
modules/cudaarithm/perf/perf_element_operations.cpp
View file @ a345402d
......@@ -1346,6 +1346,7 @@ PERF_TEST_P(Sz, MagnitudeSqr,
// Phase
DEF_PARAM_TEST(Sz_AngleInDegrees, cv::Size, bool);
DEF_PARAM_TEST(Sz_Type_AngleInDegrees, cv::Size, MatType, bool);
PERF_TEST_P(Sz_AngleInDegrees, Phase,
Combine(CUDA_TYPICAL_MAT_SIZES,
......@@ -1423,17 +1424,19 @@ PERF_TEST_P(Sz_AngleInDegrees, CartToPolar,
//////////////////////////////////////////////////////////////////////
// PolarToCart
PERF_TEST_P(Sz_AngleInDegrees, PolarToCart,
PERF_TEST_P(Sz_Type_AngleInDegrees, PolarToCart,
Combine(CUDA_TYPICAL_MAT_SIZES,
testing::Values(CV_32FC1, CV_64FC1),
Bool()))
{
const cv::Size size = GET_PARAM(0);
const bool angleInDegrees = GET_PARAM(1);
const int type = GET_PARAM(1);
const bool angleInDegrees = GET_PARAM(2);
cv::Mat magnitude(size, CV_32FC1);
cv::Mat magnitude(size, type);
declare.in(magnitude, WARMUP_RNG);
cv::Mat angle(size, CV_32FC1);
cv::Mat angle(size, type);
declare.in(angle, WARMUP_RNG);
if (PERF_RUN_CUDA())
......
modules/cudaarithm/src/cuda/polar_cart.cu
View file @ a345402d
......@@ -157,8 +157,23 @@ void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, Outpu
namespace
{
template <bool useMag>
__global__ void polarToCartImpl(const GlobPtr<float> mag, const GlobPtr<float> angle, GlobPtr<float> xmat, GlobPtr<float> ymat, const float scale, const int rows, const int cols)
template <typename T> struct sincos_op
{
__device__ __forceinline__ void operator()(T a, T *sptr, T *cptr) const
{
::sincos(a, sptr, cptr);
}
};
template <> struct sincos_op<float>
{
__device__ __forceinline__ void operator()(float a, float *sptr, float *cptr) const
{
::sincosf(a, sptr, cptr);
}
};
template <typename T, bool useMag>
__global__ void polarToCartImpl_(const GlobPtr<T> mag, const GlobPtr<T> angle, GlobPtr<T> xmat, GlobPtr<T> ymat, const T scale, const int rows, const int cols)
{
const int x = blockDim.x * blockIdx.x + threadIdx.x;
const int y = blockDim.y * blockIdx.y + threadIdx.y;
......@@ -166,45 +181,53 @@ namespace
if (x >= cols || y >= rows)
return;
const float mag_val = useMag ? mag(y, x) : 1.0f;
const float angle_val = angle(y, x);
const T mag_val = useMag ? mag(y, x) : static_cast<T>(1.0);
const T angle_val = angle(y, x);
float sin_a, cos_a;
::sincosf(scale * angle_val, &sin_a, &cos_a);
T sin_a, cos_a;
sincos_op<T> op;
op(scale * angle_val, &sin_a, &cos_a);
xmat(y, x) = mag_val * cos_a;
ymat(y, x) = mag_val * sin_a;
}
template <typename T>
void polarToCartImpl(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t& stream)
{
GpuMat_<T> xc(x.reshape(1));
GpuMat_<T> yc(y.reshape(1));
GpuMat_<T> magc(mag.reshape(1));
GpuMat_<T> anglec(angle.reshape(1));
const dim3 block(32, 8);
const dim3 grid(divUp(anglec.cols, block.x), divUp(anglec.rows, block.y));
const T scale = angleInDegrees ? static_cast<T>(CV_PI / 180.0) : static_cast<T>(1.0);
if (magc.empty())
polarToCartImpl_<T, false> << <grid, block, 0, stream >> >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
else
polarToCartImpl_<T, true> << <grid, block, 0, stream >> >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
}
}
void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _x, OutputArray _y, bool angleInDegrees, Stream& _stream)
{
typedef void(*func_t)(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t& stream);
static const func_t funcs[7] = { 0, 0, 0, 0, 0, polarToCartImpl<float>, polarToCartImpl<double> };
GpuMat mag = getInputMat(_mag, _stream);
GpuMat angle = getInputMat(_angle, _stream);
CV_Assert( angle.depth() == CV_32F );
CV_Assert(angle.depth() == CV_32F || angle.depth() == CV_64F);
CV_Assert( mag.empty() || (mag.type() == angle.type() && mag.size() == angle.size()) );
GpuMat x = getOutputMat(_x, angle.size(), CV_32FC1, _stream);
GpuMat y = getOutputMat(_y, angle.size(), CV_32FC1, _stream);
GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
GpuMat_<float> magc(mag.reshape(1));
GpuMat_<float> anglec(angle.reshape(1));
const dim3 block(32, 8);
const dim3 grid(divUp(anglec.cols, block.x), divUp(anglec.rows, block.y));
const float scale = angleInDegrees ? (CV_PI_F / 180.0f) : 1.0f;
GpuMat x = getOutputMat(_x, angle.size(), CV_MAKETYPE(angle.depth(), 1), _stream);
GpuMat y = getOutputMat(_y, angle.size(), CV_MAKETYPE(angle.depth(), 1), _stream);
cudaStream_t stream = StreamAccessor::getStream(_stream);
if (magc.empty())
polarToCartImpl<false><<<grid, block, 0, stream>>>(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
else
polarToCartImpl<true><<<grid, block, 0, stream>>>(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
funcs[angle.depth()](mag, angle, x, y, angleInDegrees, stream);
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
syncOutput(x, _x, _stream);
......
modules/cudaarithm/test/test_element_operations.cpp
View file @ a345402d
......@@ -2754,10 +2754,11 @@ INSTANTIATE_TEST_CASE_P(CUDA_Arithm, CartToPolar, testing::Combine(
////////////////////////////////////////////////////////////////////////////////
// polarToCart
PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, UseRoi)
PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, MatType, AngleInDegrees, UseRoi)
{
cv::cuda::DeviceInfo devInfo;
cv::Size size;
int type;
bool angleInDegrees;
bool useRoi;
......@@ -2765,8 +2766,9 @@ PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, Use
{
devInfo = GET_PARAM(0);
size = GET_PARAM(1);
angleInDegrees = GET_PARAM(2);
useRoi = GET_PARAM(3);
type = GET_PARAM(2);
angleInDegrees = GET_PARAM(3);
useRoi = GET_PARAM(4);
cv::cuda::setDevice(devInfo.deviceID());
}
......@@ -2774,24 +2776,26 @@ PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, Use
CUDA_TEST_P(PolarToCart, Accuracy)
{
cv::Mat magnitude = randomMat(size, CV_32FC1);
cv::Mat angle = randomMat(size, CV_32FC1);
cv::Mat magnitude = randomMat(size, type);
cv::Mat angle = randomMat(size, type);
const double tol = (type == CV_32FC1 ? 1.6e-4 : 1e-4) * (angleInDegrees ? 1.0 : 19.0);
cv::cuda::GpuMat x = createMat(size, CV_32FC1, useRoi);
cv::cuda::GpuMat y = createMat(size, CV_32FC1, useRoi);
cv::cuda::GpuMat x = createMat(size, type, useRoi);
cv::cuda::GpuMat y = createMat(size, type, useRoi);
cv::cuda::polarToCart(loadMat(magnitude, useRoi), loadMat(angle, useRoi), x, y, angleInDegrees);
cv::Mat x_gold;
cv::Mat y_gold;
cv::polarToCart(magnitude, angle, x_gold, y_gold, angleInDegrees);
EXPECT_MAT_NEAR(x_gold, x, 1e-4);
EXPECT_MAT_NEAR(y_gold, y, 1e-4);
EXPECT_MAT_NEAR(x_gold, x, tol);
EXPECT_MAT_NEAR(y_gold, y, tol);
}
INSTANTIATE_TEST_CASE_P(CUDA_Arithm, PolarToCart, testing::Combine(
ALL_DEVICES,
DIFFERENT_SIZES,
testing::Values(CV_32FC1, CV_64FC1),
testing::Values(AngleInDegrees(false), AngleInDegrees(true)),
WHOLE_SUBMAT));
......
modules/cudaimgproc/test/test_color.cpp
View file @ a345402d
......@@ -1740,7 +1740,7 @@ CUDA_TEST_P(CvtColor, Lab2BGR)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Lab2BGR);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-5);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-5);
}
CUDA_TEST_P(CvtColor, Lab2RGB)
......@@ -1757,7 +1757,7 @@ CUDA_TEST_P(CvtColor, Lab2RGB)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Lab2RGB);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-5);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-5);
}
CUDA_TEST_P(CvtColor, Lab2BGRA)
......@@ -1776,7 +1776,7 @@ CUDA_TEST_P(CvtColor, Lab2BGRA)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Lab2BGR, 4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-5);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-5);
}
CUDA_TEST_P(CvtColor, Lab2LBGR)
......@@ -1793,7 +1793,7 @@ CUDA_TEST_P(CvtColor, Lab2LBGR)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Lab2LBGR);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-5);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-5);
}
CUDA_TEST_P(CvtColor, Lab2LRGB)
......@@ -1810,7 +1810,7 @@ CUDA_TEST_P(CvtColor, Lab2LRGB)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Lab2LRGB);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-5);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-5);
}
CUDA_TEST_P(CvtColor, Lab2LRGBA)
......@@ -1827,7 +1827,7 @@ CUDA_TEST_P(CvtColor, Lab2LRGBA)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Lab2LRGB, 4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-5);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-5);
}
CUDA_TEST_P(CvtColor, BGR2Luv)
......@@ -1958,7 +1958,7 @@ CUDA_TEST_P(CvtColor, Luv2BGR)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Luv2BGR);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-4);
}
CUDA_TEST_P(CvtColor, Luv2RGB)
......@@ -1975,7 +1975,7 @@ CUDA_TEST_P(CvtColor, Luv2RGB)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Luv2RGB);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-4);
}
CUDA_TEST_P(CvtColor, Luv2BGRA)
......@@ -1994,7 +1994,7 @@ CUDA_TEST_P(CvtColor, Luv2BGRA)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Luv2BGR, 4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-4);
}
CUDA_TEST_P(CvtColor, Luv2LBGR)
......@@ -2011,7 +2011,7 @@ CUDA_TEST_P(CvtColor, Luv2LBGR)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Luv2LBGR);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-4);
}
CUDA_TEST_P(CvtColor, Luv2LRGB)
......@@ -2028,7 +2028,7 @@ CUDA_TEST_P(CvtColor, Luv2LRGB)
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_Luv2LRGB);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-4);
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 2 : 1e-4);
}
CUDA_TEST_P(CvtColor, Luv2LRGBA)
......
modules/cudev/include/opencv2/cudev/functional/detail/color_cvt.hpp
View file @ a345402d
......@@ -1207,27 +1207,29 @@ namespace color_cvt_detail
__device__ typename MakeVec<float, dcn>::type operator ()(const typename MakeVec<float, scn>::type& src) const
{
const float _d = 1.f / (0.950456f + 15 + 1.088754f * 3);
const float _un = 4 * 0.950456f * _d;
const float _vn = 9 * _d;
const float _un = 13 * 4 * 0.950456f * _d;
const float _vn = 13 * 9 * _d;
float L = src.x;
float u = src.y;
float v = src.z;
float Y = (L + 16.f) * (1.f / 116.f);
Y = Y * Y * Y;
float Y1 = (L + 16.f) * (1.f / 116.f);
Y1 = Y1 * Y1 * Y1;
float Y0 = L * (1.f / 903.3f);
float Y = L <= 8.f ? Y0 : Y1;
float d = (1.f / 13.f) / L;
u = u * d + _un;
v = v * d + _vn;
u = (u + _un * L) * 3.f;
v = (v + _vn * L) * 4.f;
float iv = 1.f / v;
float X = 2.25f * u * Y * iv;
float Z = (12 - 3 * u - 20 * v) * Y * 0.25f * iv;
iv = ::fmaxf(-0.25f, ::fminf(0.25f, iv));
float X = 3.f * u * iv;
float Z = (12.f * 13.f * L - u) * iv - 5.f;
float B = 0.055648f * X - 0.204043f * Y + 1.057311f * Z;
float G = -0.969256f * X + 1.875991f * Y + 0.041556f * Z;
float R = 3.240479f * X - 1.537150f * Y - 0.498535f * Z;
float B = (0.055648f * X - 0.204043f + 1.057311f * Z) * Y;
float G = (-0.969256f * X + 1.875991f + 0.041556f * Z) * Y;
float R = (3.240479f * X - 1.537150f - 0.498535f * Z) * Y;
R = ::fminf(::fmaxf(R, 0.f), 1.f);
G = ::fminf(::fmaxf(G, 0.f), 1.f);
......
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