{
float* btvWeights_ = NULL;
size_t btvWeights_size = 0;
+ oclMat c_btvRegWeights;
}
}
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);
}
-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();
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));
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
{
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);
}
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;
oclMat highRes_;
vector<oclMat> diffTerms_;
- vector<oclMat> a_, b_, c_;
+ oclMat a_, b_, c_, d_;
oclMat regTerm_;
};
btvKernelSize_ = 7;
blurKernelSize_ = 5;
blurSigma_ = 0.0;
- opticalFlow_ = createOptFlow_DualTVL1_OCL();
+ opticalFlow_ = createOptFlow_Farneback_OCL();
curBlurKernelSize_ = -1;
curBlurSigma_ = -1.0;
// 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)
highRes_.release();
diffTerms_.clear();
- a_.clear();
- b_.clear();
- c_.clear();
+ a_.release();
+ b_.release();
+ c_.release();
regTerm_.release();
+ c_btvRegWeights.release();
}
////////////////////////////////////////////////////////////
//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);
}
__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);
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);
}
}
}
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;
}
__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);
{
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;
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)
{
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);
}
}
}
projects / profile / ivi / opencv.git / commitdiff