generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

debug of cuda_tvl1 => pass tests succesfully

Browse Source
This commit is contained in:
Ernest Galbrun 2014-07-04 14:23:09 +02:00
parent 32707317fa
commit 693c4e5741
4 changed files with 34 additions and 23 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
modules/cudaoptflow/src/cuda/tvl1flow.cu
View File

@ -209,7 +209,9 @@ namespace tvl1flow
__global__ void estimateUKernel(const PtrStepSzf I1wx, const PtrStepf I1wy, __global__ void estimateUKernel(const PtrStepSzf I1wx, const PtrStepf I1wy,
const PtrStepf grad, const PtrStepf rho_c, const PtrStepf grad, const PtrStepf rho_c,
const PtrStepf p11, const PtrStepf p12, const PtrStepf p21, const PtrStepf p22, const PtrStepf p31, const PtrStepf p32, const PtrStepf p11, const PtrStepf p12,
const PtrStepf p21, const PtrStepf p22,
const PtrStepf p31, const PtrStepf p32,
PtrStepf u1, PtrStepf u2, PtrStepf u3, PtrStepf error, PtrStepf u1, PtrStepf u2, PtrStepf u3, PtrStepf error,
const float l_t, const float theta, const float gamma, const bool calcError) const float l_t, const float theta, const float gamma, const bool calcError)
{ {

12
modules/cudaoptflow/src/tvl1flow.cpp
View File

@ -235,9 +235,8 @@ void cv::cuda::OpticalFlowDual_TVL1_CUDA::procOneScale(const GpuMat& I0, const G
{ {
// some tweaks to make sum operation less frequently // some tweaks to make sum operation less frequently
bool calcError = (epsilon > 0) && (n & 0x1) && (prevError < scaledEpsilon); bool calcError = (epsilon > 0) && (n & 0x1) && (prevError < scaledEpsilon);
cv::Mat m1(u3);
estimateU(I1wx, I1wy, grad, rho_c, p11, p12, p21, p22, p31, p32, u1, u2, u3, diff, l_t, gamma, static_cast<float>(theta), calcError); estimateU(I1wx, I1wy, grad, rho_c, p11, p12, p21, p22, p31, p32, u1, u2, u3, diff, l_t, static_cast<float>(theta), gamma, calcError);
if (calcError) if (calcError)
{ {
error = cuda::sum(diff, norm_buf)[0]; error = cuda::sum(diff, norm_buf)[0];
@ -259,7 +258,8 @@ void cv::cuda::OpticalFlowDual_TVL1_CUDA::collectGarbage()
I0s.clear(); I0s.clear();
I1s.clear(); I1s.clear();
u1s.clear(); u1s.clear();
u2s.clear(); u2s.clear();
u3s.clear();
I1x_buf.release(); I1x_buf.release();
I1y_buf.release(); I1y_buf.release();
@ -274,7 +274,9 @@ void cv::cuda::OpticalFlowDual_TVL1_CUDA::collectGarbage()
p11_buf.release(); p11_buf.release();
p12_buf.release(); p12_buf.release();
p21_buf.release(); p21_buf.release();
p22_buf.release(); p22_buf.release();
p31_buf.release();
p32_buf.release();
diff_buf.release(); diff_buf.release();
norm_buf.release(); norm_buf.release();

12
modules/cudaoptflow/test/test_optflow.cpp
View File

@ -361,9 +361,21 @@ CUDA_TEST_P(OpticalFlowDual_TVL1, Accuracy)
alg->calc(frame0, frame1, flow); alg->calc(frame0, frame1, flow);
cv::Mat gold[2]; cv::Mat gold[2];
cv::split(flow, gold); cv::split(flow, gold);
cv::Mat mx(d_flowx);
cv::Mat my(d_flowx);
EXPECT_MAT_SIMILAR(gold[0], d_flowx, 4e-3); EXPECT_MAT_SIMILAR(gold[0], d_flowx, 4e-3);
EXPECT_MAT_SIMILAR(gold[1], d_flowy, 4e-3); EXPECT_MAT_SIMILAR(gold[1], d_flowy, 4e-3);
d_alg.gamma = 1;
alg->set("gamma", 1);
d_alg(loadMat(frame0, useRoi), loadMat(frame1, useRoi), d_flowx, d_flowy);
alg->calc(frame0, frame1, flow);
cv::split(flow, gold);
mx = cv::Mat(d_flowx);
my = cv::Mat(d_flowx);
EXPECT_MAT_SIMILAR(gold[0], d_flowx, 4e-3);
EXPECT_MAT_SIMILAR(gold[1], d_flowy, 4e-3);
} }
INSTANTIATE_TEST_CASE_P(CUDA_OptFlow, OpticalFlowDual_TVL1, testing::Combine( INSTANTIATE_TEST_CASE_P(CUDA_OptFlow, OpticalFlowDual_TVL1, testing::Combine(

29
modules/video/src/tvl1flow.cpp
View File

@ -121,8 +121,8 @@ private:
std::vector<Mat_<float> > I0s; std::vector<Mat_<float> > I0s;
std::vector<Mat_<float> > I1s; std::vector<Mat_<float> > I1s;
std::vector<Mat_<float> > u1s; std::vector<Mat_<float> > u1s;
std::vector<Mat_<float> > u2s; std::vector<Mat_<float> > u2s;
std::vector<Mat_<float> > u3s; std::vector<Mat_<float> > u3s;
Mat_<float> I1x_buf; Mat_<float> I1x_buf;
Mat_<float> I1y_buf; Mat_<float> I1y_buf;
@ -138,26 +138,26 @@ private:
Mat_<float> rho_c_buf; Mat_<float> rho_c_buf;
Mat_<float> v1_buf; Mat_<float> v1_buf;
Mat_<float> v2_buf; Mat_<float> v2_buf;
Mat_<float> v3_buf; Mat_<float> v3_buf;
Mat_<float> p11_buf; Mat_<float> p11_buf;
Mat_<float> p12_buf; Mat_<float> p12_buf;
Mat_<float> p21_buf; Mat_<float> p21_buf;
Mat_<float> p22_buf; Mat_<float> p22_buf;
Mat_<float> p31_buf; Mat_<float> p31_buf;
Mat_<float> p32_buf; Mat_<float> p32_buf;
Mat_<float> div_p1_buf; Mat_<float> div_p1_buf;
Mat_<float> div_p2_buf; Mat_<float> div_p2_buf;
Mat_<float> div_p3_buf; Mat_<float> div_p3_buf;
Mat_<float> u1x_buf; Mat_<float> u1x_buf;
Mat_<float> u1y_buf; Mat_<float> u1y_buf;
Mat_<float> u2x_buf; Mat_<float> u2x_buf;
Mat_<float> u2y_buf; Mat_<float> u2y_buf;
Mat_<float> u3x_buf; Mat_<float> u3x_buf;
Mat_<float> u3y_buf; Mat_<float> u3y_buf;
} dm; } dm;
struct dataUMat struct dataUMat
{ {
@ -892,10 +892,6 @@ void CalcGradRhoBody::operator() (const Range& range) const
// compute the constant part of the rho function // compute the constant part of the rho function
rhoRow[x] = (I1wRow[x] - I1wxRow[x] * u1Row[x] - I1wyRow[x] * u2Row[x] - I0Row[x]); rhoRow[x] = (I1wRow[x] - I1wxRow[x] * u1Row[x] - I1wyRow[x] * u2Row[x] - I0Row[x]);
//It = I1wRow[x] - I0Row[x]
//(u - u0)*i_X = I1wxRow[x] * u1Row[x]
//(v - v0)*i_Y = I1wyRow[x] * u2Row[x]
// gamma * w = gamma * u3
} }
} }
} }
@ -970,7 +966,6 @@ void EstimateVBody::operator() (const Range& range) const
float d1 = 0.0f; float d1 = 0.0f;
float d2 = 0.0f; float d2 = 0.0f;
float d3 = 0.0f; float d3 = 0.0f;
// add d3 for 3 cases
if (rho < -l_t * gradRow[x]) if (rho < -l_t * gradRow[x])
{ {
d1 = l_t * I1wxRow[x]; d1 = l_t * I1wxRow[x];