Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

1. Let btvRegWeights to be constant per `process` call.

Browse Source 
2. Let Farneback to be the default optical flow method.
3. Fix a timing method bug for ocl path.
4. Remove useless finish operation in farneback
5. Refactor buffer usage.
pull/1655/head
peng xiao 12 years ago
parent 1eb322fa5d
commit 7285341083
4 changed files with 84 additions and 165 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 2
      modules/ocl/src/optical_flow_farneback.cpp
  2. 78
      modules/superres/src/btv_l1_ocl.cpp
  3. 163
      modules/superres/src/opencl/superres_btvl1.cl
  4. 6
      samples/gpu/super_resolution.cpp

2
modules/ocl/src/optical_flow_farneback.cpp
Unescape View File

@ -336,8 +336,6 @@ void cv::ocl::FarnebackOpticalFlow::updateFlow_boxFilter(
swap(M, bufM);
finish();
optflow_farneback::updateFlowOcl(M, flowx, flowy);
if (updateMatrices)

78
modules/superres/src/btv_l1_ocl.cpp
Unescape View File

@ -70,6 +70,7 @@ namespace cv
{
float* btvWeights_ = NULL;
size_t btvWeights_size = 0;
oclMat c_btvRegWeights;
}
}
@ -82,10 +83,6 @@ namespace btv_l1_device_ocl
void upscale(const oclMat& src, oclMat& dst, int scale);
float diffSign(float a, float b);
Point3f diffSign(Point3f a, Point3f b);
void diffSign(const oclMat& src1, const oclMat& src2, oclMat& dst);
void calcBtvRegularization(const oclMat& src, oclMat& dst, int ksize);
@ -165,20 +162,6 @@ void btv_l1_device_ocl::upscale(const oclMat& src, oclMat& dst, int scale)
}
float btv_l1_device_ocl::diffSign(float a, float b)
{
return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
}
Point3f btv_l1_device_ocl::diffSign(Point3f a, Point3f b)
{
return Point3f(
a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f,
a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f,
a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f
);
}
void btv_l1_device_ocl::diffSign(const oclMat& src1, const oclMat& src2, oclMat& dst)
{
Context* clCxt = Context::getContext();
@ -228,12 +211,6 @@ void btv_l1_device_ocl::calcBtvRegularization(const oclMat& src, oclMat& dst, in
int cn = src.oclchannels();
cl_mem c_btvRegWeights;
size_t count = btvWeights_size * sizeof(float);
c_btvRegWeights = openCLCreateBuffer(clCxt, CL_MEM_READ_ONLY, count);
int cl_safe_check = clEnqueueWriteBuffer(getClCommandQueue(clCxt), c_btvRegWeights, 1, 0, count, btvWeights_, 0, NULL, NULL);
CV_Assert(cl_safe_check == CL_SUCCESS);
args.push_back(make_pair(sizeof(cl_mem), (void*)&src_.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&dst_.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&src_step));
@ -242,11 +219,9 @@ void btv_l1_device_ocl::calcBtvRegularization(const oclMat& src, oclMat& dst, in
args.push_back(make_pair(sizeof(cl_int), (void*)&src.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&ksize));
args.push_back(make_pair(sizeof(cl_int), (void*)&cn));
args.push_back(make_pair(sizeof(cl_mem), (void*)&c_btvRegWeights));
args.push_back(make_pair(sizeof(cl_mem), (void*)&c_btvRegWeights.data));
openCLExecuteKernel(clCxt, &superres_btvl1, kernel_name, global_thread, local_thread, args, -1, -1);
cl_safe_check = clReleaseMemObject(c_btvRegWeights);
CV_Assert(cl_safe_check == CL_SUCCESS);
}
namespace
@ -321,9 +296,6 @@ namespace
{
CV_Assert( src.channels() == 1 || src.channels() == 3 || src.channels() == 4 );
dst.create(src.rows * scale, src.cols * scale, src.type());
dst.setTo(Scalar::all(0));
btv_l1_device_ocl::upscale(src, dst, scale);
}
@ -351,12 +323,13 @@ namespace
btvWeights_ = &btvWeights[0];
btvWeights_size = size;
Mat btvWeights_mheader(1, static_cast<int>(size), CV_32FC1, btvWeights_);
c_btvRegWeights = btvWeights_mheader;
}
void calcBtvRegularization(const oclMat& src, oclMat& dst, int btvKernelSize)
{
dst.create(src.size(), src.type());
dst.setTo(Scalar::all(0));
const int ksize = (btvKernelSize - 1) / 2;
@ -407,7 +380,7 @@ namespace
oclMat highRes_;
vector<oclMat> diffTerms_;
vector<oclMat> a_, b_, c_;
oclMat a_, b_, c_, d_;
oclMat regTerm_;
};
@ -421,7 +394,7 @@ namespace
btvKernelSize_ = 7;
blurKernelSize_ = 5;
blurSigma_ = 0.0;
opticalFlow_ = createOptFlow_DualTVL1_OCL();
opticalFlow_ = createOptFlow_Farneback_OCL();
curBlurKernelSize_ = -1;
curBlurSigma_ = -1.0;
@ -487,34 +460,36 @@ namespace
// iterations
diffTerms_.resize(src.size());
a_.resize(src.size());
b_.resize(src.size());
c_.resize(src.size());
bool d_inited = false;
a_.create(highRes_.size(), highRes_.type());
b_.create(highRes_.size(), highRes_.type());
c_.create(lowResSize, highRes_.type());
d_.create(highRes_.rows, highRes_.cols, highRes_.type());
for (int i = 0; i < iterations_; ++i)
{
if(!d_inited)
{
d_.setTo(0);
d_inited = true;
}
for (size_t k = 0; k < src.size(); ++k)
{
diffTerms_[k].create(highRes_.size(), highRes_.type());
a_[k].create(highRes_.size(), highRes_.type());
b_[k].create(highRes_.size(), highRes_.type());
c_[k].create(lowResSize, highRes_.type());
// a = M * Ih
ocl::remap(highRes_, a_[k], backwardMaps_[k].first, backwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar());
ocl::remap(highRes_, a_, backwardMaps_[k].first, backwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar());
// b = HM * Ih
filters_[k]->apply(a_[k], b_[k], Rect(0,0,-1,-1));
filters_[k]->apply(a_, b_, Rect(0,0,-1,-1));
// c = DHF * Ih
ocl::resize(b_[k], c_[k], lowResSize, 0, 0, INTER_NEAREST);
ocl::resize(b_, c_, lowResSize, 0, 0, INTER_NEAREST);
diffSign(src[k], c_[k], c_[k]);
diffSign(src[k], c_, c_);
// a = Dt * diff
upscale(c_[k], a_[k], scale_);
upscale(c_, d_, scale_);
// b = HtDt * diff
filters_[k]->apply(a_[k], b_[k], Rect(0,0,-1,-1));
filters_[k]->apply(d_, b_, Rect(0,0,-1,-1));
// diffTerm = MtHtDt * diff
ocl::remap(b_[k], diffTerms_[k], forwardMaps_[k].first, forwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar());
ocl::remap(b_, diffTerms_[k], forwardMaps_[k].first, forwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar());
}
if (lambda_ > 0)
@ -549,10 +524,11 @@ namespace
highRes_.release();
diffTerms_.clear();
a_.clear();
b_.clear();
c_.clear();
a_.release();
b_.release();
c_.release();
regTerm_.release();
c_btvRegWeights.release();
}
////////////////////////////////////////////////////////////

163
modules/superres/src/opencl/superres_btvl1.cl
Unescape View File

@ -44,24 +44,24 @@
//M*/
__kernel void buildMotionMapsKernel(__global float* forwardMotionX,
__global float* forwardMotionY,
__global float* backwardMotionX,
__global float* backwardMotionY,
__global float* forwardMapX,
__global float* forwardMapY,
__global float* backwardMapX,
__global float* backwardMapY,
int forwardMotionX_row,
int forwardMotionX_col,
int forwardMotionX_step,
int forwardMotionY_step,
int backwardMotionX_step,
int backwardMotionY_step,
int forwardMapX_step,
int forwardMapY_step,
int backwardMapX_step,
int backwardMapY_step
)
__global float* forwardMotionY,
__global float* backwardMotionX,
__global float* backwardMotionY,
__global float* forwardMapX,
__global float* forwardMapY,
__global float* backwardMapX,
__global float* backwardMapY,
int forwardMotionX_row,
int forwardMotionX_col,
int forwardMotionX_step,
int forwardMotionY_step,
int backwardMotionX_step,
int backwardMotionY_step,
int forwardMapX_step,
int forwardMapY_step,
int backwardMapX_step,
int backwardMapY_step
)
{
int x = get_global_id(0);
int y = get_global_id(1);
@ -83,14 +83,14 @@ __kernel void buildMotionMapsKernel(__global float* forwardMotionX,
}
__kernel void upscaleKernel(__global float* src,
__global float* dst,
int src_step,
int dst_step,
int src_row,
int src_col,
int scale,
int channels
)
__global float* dst,
int src_step,
int dst_step,
int src_row,
int src_col,
int scale,
int channels
)
{
int x = get_global_id(0);
int y = get_global_id(1);
@ -100,17 +100,10 @@ __kernel void upscaleKernel(__global float* src,
if(channels == 1)
{
dst[y * scale * dst_step + x * scale] = src[y * src_step + x];
}else if(channels == 3)
{
dst[y * channels * scale * dst_step + 3 * x * scale + 0] = src[y * channels * src_step + 3 * x + 0];
dst[y * channels * scale * dst_step + 3 * x * scale + 1] = src[y * channels * src_step + 3 * x + 1];
dst[y * channels * scale * dst_step + 3 * x * scale + 2] = src[y * channels * src_step + 3 * x + 2];
}else
}
else
{
dst[y * channels * scale * dst_step + 4 * x * scale + 0] = src[y * channels * src_step + 4 * x + 0];
dst[y * channels * scale * dst_step + 4 * x * scale + 1] = src[y * channels * src_step + 4 * x + 1];
dst[y * channels * scale * dst_step + 4 * x * scale + 2] = src[y * channels * src_step + 4 * x + 2];
dst[y * channels * scale * dst_step + 4 * x * scale + 3] = src[y * channels * src_step + 4 * x + 3];
vstore4(vload4(0, src + y * channels * src_step + 4 * x), 0, dst + y * channels * scale * dst_step + 4 * x * scale);
}
}
}
@ -121,15 +114,6 @@ float diffSign(float a, float b)
return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
}
float3 diffSign3(float3 a, float3 b)
{
float3 pos;
pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
return pos;
}
float4 diffSign4(float4 a, float4 b)
{
float4 pos;
@ -141,13 +125,13 @@ float4 diffSign4(float4 a, float4 b)
}
__kernel void diffSignKernel(__global float* src1,
__global float* src2,
__global float* dst,
int src1_row,
int src1_col,
int dst_step,
int src1_step,
int src2_step)
__global float* src2,
__global float* dst,
int src1_row,
int src1_col,
int dst_step,
int src1_step,
int src2_step)
{
int x = get_global_id(0);
int y = get_global_id(1);
@ -156,19 +140,18 @@ __kernel void diffSignKernel(__global float* src1,
{
dst[y * dst_step + x] = diffSign(src1[y * src1_step + x], src2[y * src2_step + x]);
}
barrier(CLK_LOCAL_MEM_FENCE);
}
__kernel void calcBtvRegularizationKernel(__global float* src,
__global float* dst,
int src_step,
int dst_step,
int src_row,
int src_col,
int ksize,
int channels,
__global float* c_btvRegWeights
)
__global float* dst,
int src_step,
int dst_step,
int src_row,
int src_col,
int ksize,
int channels,
__constant float* c_btvRegWeights
)
{
int x = get_global_id(0) + ksize;
int y = get_global_id(1) + ksize;
@ -180,57 +163,19 @@ __kernel void calcBtvRegularizationKernel(__global float* src,
const float srcVal = src[y * src_step + x];
float dstVal = 0.0f;
for (int m = 0, count = 0; m <= ksize; ++m)
{
for (int l = ksize; l + m >= 0; --l, ++count)
dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal));
}
dst[y * dst_step + x] = dstVal;
}else if(channels == 3)
{
float3 srcVal;
srcVal.x = src[y * src_step + 3 * x + 0];
srcVal.y = src[y * src_step + 3 * x + 1];
srcVal.z = src[y * src_step + 3 * x + 2];
float3 dstVal;
dstVal.x = 0.0f;
dstVal.y = 0.0f;
dstVal.z = 0.0f;
for (int m = 0, count = 0; m <= ksize; ++m)
{
for (int l = ksize; l + m >= 0; --l, ++count)
{
float3 src1;
src1.x = src[(y + m) * src_step + 3 * (x + l) + 0];
src1.y = src[(y + m) * src_step + 3 * (x + l) + 1];
src1.z = src[(y + m) * src_step + 3 * (x + l) + 2];
float3 src2;
src2.x = src[(y - m) * src_step + 3 * (x - l) + 0];
src2.y = src[(y - m) * src_step + 3 * (x - l) + 1];
src2.z = src[(y - m) * src_step + 3 * (x - l) + 2];
dstVal = dstVal + c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal));
dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal));
}
}
dst[y * dst_step + 3 * x + 0] = dstVal.x;
dst[y * dst_step + 3 * x + 1] = dstVal.y;
dst[y * dst_step + 3 * x + 2] = dstVal.z;
}else
dst[y * dst_step + x] = dstVal;
}
else
{
float4 srcVal;
srcVal.x = src[y * src_step + 4 * x + 0];//r type =float
srcVal.y = src[y * src_step + 4 * x + 1];//g
srcVal.z = src[y * src_step + 4 * x + 2];//b
srcVal.w = src[y * src_step + 4 * x + 3];//a
float4 dstVal;
dstVal.x = 0.0f;
dstVal.y = 0.0f;
dstVal.z = 0.0f;
dstVal.w = 0.0f;
float4 srcVal = vload4(0, src + y * src_step + 4 * x);
float4 dstVal = 0.f;
for (int m = 0, count = 0; m <= ksize; ++m)
{
@ -249,13 +194,9 @@ __kernel void calcBtvRegularizationKernel(__global float* src,
src2.w = src[(y - m) * src_step + 4 * (x - l) + 3];
dstVal = dstVal + c_btvRegWeights[count] * (diffSign4(srcVal, src1) - diffSign4(src2, srcVal));
}
}
dst[y * dst_step + 4 * x + 0] = dstVal.x;
dst[y * dst_step + 4 * x + 1] = dstVal.y;
dst[y * dst_step + 4 * x + 2] = dstVal.z;
dst[y * dst_step + 4 * x + 3] = dstVal.w;
vstore4(dstVal, 0, dst + y * dst_step + 4 * x);
}
}
}

6
samples/gpu/super_resolution.cpp
Unescape View File

@ -221,7 +221,11 @@ int main(int argc, const char* argv[])
if(useOcl)
{
MEASURE_TIME(superRes->nextFrame(result_));
MEASURE_TIME(
{
superRes->nextFrame(result_);
ocl::finish();
});
}
else
#endif

Write Preview
Loading…
Cancel
Save