Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

added ROI support to ocl::buildWarp*Maps functions

Browse Source
pull/1773/head
Ilya Lavrenov 11 years ago
parent 0be2752366
commit 4248f82211
2 changed files with 142 additions and 153 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. 146
      modules/ocl/src/build_warps.cpp
  2. 149
      modules/ocl/src/opencl/build_warps.cl

146
modules/ocl/src/build_warps.cpp
Unescape View File

@ -53,7 +53,7 @@ using namespace cv::ocl;
// buildWarpPlaneMaps
void cv::ocl::buildWarpPlaneMaps(Size /*src_size*/, Rect dst_roi, const Mat &K, const Mat &R, const Mat &T,
float scale, oclMat &map_x, oclMat &map_y)
float scale, oclMat &xmap, oclMat &ymap)
{
CV_Assert(K.size() == Size(3, 3) && K.type() == CV_32F);
CV_Assert(R.size() == Size(3, 3) && R.type() == CV_32F);
@ -68,37 +68,40 @@ void cv::ocl::buildWarpPlaneMaps(Size /*src_size*/, Rect dst_roi, const Mat &K,
oclMat KRT_oclMat(KRT_mat);
// transfer K_Rinv and T into a single cl_mem
map_x.create(dst_roi.size(), CV_32F);
map_y.create(dst_roi.size(), CV_32F);
xmap.create(dst_roi.size(), CV_32F);
ymap.create(dst_roi.size(), CV_32F);
int tl_u = dst_roi.tl().x;
int tl_v = dst_roi.tl().y;
Context *clCxt = Context::getContext();
string kernelName = "buildWarpPlaneMaps";
vector< pair<size_t, const void *> > args;
int xmap_step = xmap.step / xmap.elemSize(), xmap_offset = xmap.offset / xmap.elemSize();
int ymap_step = ymap.step / ymap.elemSize(), ymap_offset = ymap.offset / ymap.elemSize();
args.push_back( make_pair( sizeof(cl_mem), (void *)&map_x.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&map_y.data));
vector< pair<size_t, const void *> > args;
args.push_back( make_pair( sizeof(cl_mem), (void *)&xmap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&ymap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&KRT_mat.data));
args.push_back( make_pair( sizeof(cl_int), (void *)&tl_u));
args.push_back( make_pair( sizeof(cl_int), (void *)&tl_v));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_y.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_offset));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_offset));
args.push_back( make_pair( sizeof(cl_float), (void *)&scale));
size_t globalThreads[3] = {map_x.cols, map_x.rows, 1};
size_t localThreads[3] = {32, 8, 1};
openCLExecuteKernel(clCxt, &build_warps, kernelName, globalThreads, localThreads, args, -1, -1);
size_t globalThreads[3] = { xmap.cols, xmap.rows, 1 };
size_t localThreads[3] = { 32, 8, 1 };
openCLExecuteKernel(Context::getContext(), &build_warps, "buildWarpPlaneMaps", globalThreads, localThreads, args, -1, -1);
}
//////////////////////////////////////////////////////////////////////////////
// buildWarpCylyndricalMaps
void cv::ocl::buildWarpCylindricalMaps(Size /*src_size*/, Rect dst_roi, const Mat &K, const Mat &R, float scale,
oclMat &map_x, oclMat &map_y)
oclMat &xmap, oclMat &ymap)
{
CV_Assert(K.size() == Size(3, 3) && K.type() == CV_32F);
CV_Assert(R.size() == Size(3, 3) && R.type() == CV_32F);
@ -108,36 +111,40 @@ void cv::ocl::buildWarpCylindricalMaps(Size /*src_size*/, Rect dst_roi, const Ma
oclMat KR_oclMat(K_Rinv.reshape(1, 1));
map_x.create(dst_roi.size(), CV_32F);
map_y.create(dst_roi.size(), CV_32F);
xmap.create(dst_roi.size(), CV_32F);
ymap.create(dst_roi.size(), CV_32F);
int tl_u = dst_roi.tl().x;
int tl_v = dst_roi.tl().y;
Context *clCxt = Context::getContext();
string kernelName = "buildWarpCylindricalMaps";
vector< pair<size_t, const void *> > args;
int xmap_step = xmap.step / xmap.elemSize(), xmap_offset = xmap.offset / xmap.elemSize();
int ymap_step = ymap.step / ymap.elemSize(), ymap_offset = ymap.offset / ymap.elemSize();
args.push_back( make_pair( sizeof(cl_mem), (void *)&map_x.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&map_y.data));
vector< pair<size_t, const void *> > args;
args.push_back( make_pair( sizeof(cl_mem), (void *)&xmap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&ymap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&KR_oclMat.data));
args.push_back( make_pair( sizeof(cl_int), (void *)&tl_u));
args.push_back( make_pair( sizeof(cl_int), (void *)&tl_v));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_y.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_offset));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_offset));
args.push_back( make_pair( sizeof(cl_float), (void *)&scale));
size_t globalThreads[3] = {map_x.cols, map_x.rows, 1};
size_t localThreads[3] = {32, 8, 1};
openCLExecuteKernel(clCxt, &build_warps, kernelName, globalThreads, localThreads, args, -1, -1);
size_t globalThreads[3] = { xmap.cols, xmap.rows, 1 };
size_t localThreads[3] = { 32, 8, 1 };
openCLExecuteKernel(Context::getContext(), &build_warps, "buildWarpCylindricalMaps", globalThreads, localThreads, args, -1, -1);
}
//////////////////////////////////////////////////////////////////////////////
// buildWarpSphericalMaps
void cv::ocl::buildWarpSphericalMaps(Size /*src_size*/, Rect dst_roi, const Mat &K, const Mat &R, float scale,
oclMat &map_x, oclMat &map_y)
oclMat &xmap, oclMat &ymap)
{
CV_Assert(K.size() == Size(3, 3) && K.type() == CV_32F);
CV_Assert(R.size() == Size(3, 3) && R.type() == CV_32F);
@ -147,37 +154,41 @@ void cv::ocl::buildWarpSphericalMaps(Size /*src_size*/, Rect dst_roi, const Mat
oclMat KR_oclMat(K_Rinv.reshape(1, 1));
// transfer K_Rinv, R_Kinv into a single cl_mem
map_x.create(dst_roi.size(), CV_32F);
map_y.create(dst_roi.size(), CV_32F);
xmap.create(dst_roi.size(), CV_32F);
ymap.create(dst_roi.size(), CV_32F);
int tl_u = dst_roi.tl().x;
int tl_v = dst_roi.tl().y;
Context *clCxt = Context::getContext();
string kernelName = "buildWarpSphericalMaps";
vector< pair<size_t, const void *> > args;
int xmap_step = xmap.step / xmap.elemSize(), xmap_offset = xmap.offset / xmap.elemSize();
int ymap_step = ymap.step / ymap.elemSize(), ymap_offset = ymap.offset / ymap.elemSize();
args.push_back( make_pair( sizeof(cl_mem), (void *)&map_x.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&map_y.data));
vector< pair<size_t, const void *> > args;
args.push_back( make_pair( sizeof(cl_mem), (void *)&xmap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&ymap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&KR_oclMat.data));
args.push_back( make_pair( sizeof(cl_int), (void *)&tl_u));
args.push_back( make_pair( sizeof(cl_int), (void *)&tl_v));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_x.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&map_y.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_offset));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_offset));
args.push_back( make_pair( sizeof(cl_float), (void *)&scale));
size_t globalThreads[3] = {map_x.cols, map_x.rows, 1};
size_t localThreads[3] = {32, 8, 1};
openCLExecuteKernel(clCxt, &build_warps, kernelName, globalThreads, localThreads, args, -1, -1);
size_t globalThreads[3] = { xmap.cols, xmap.rows, 1 };
size_t localThreads[3] = { 32, 8, 1 };
openCLExecuteKernel(Context::getContext(), &build_warps, "buildWarpSphericalMaps", globalThreads, localThreads, args, -1, -1);
}
//////////////////////////////////////////////////////////////////////////////
// buildWarpAffineMaps
void cv::ocl::buildWarpAffineMaps(const Mat &M, bool inverse, Size dsize, oclMat &xmap, oclMat &ymap)
{
CV_Assert(M.rows == 2 && M.cols == 3);
CV_Assert(dsize.area());
xmap.create(dsize, CV_32FC1);
ymap.create(dsize, CV_32FC1);
@ -194,29 +205,34 @@ void cv::ocl::buildWarpAffineMaps(const Mat &M, bool inverse, Size dsize, oclMat
iM.convertTo(coeffsMat, coeffsMat.type());
}
int xmap_step = xmap.step / xmap.elemSize(), xmap_offset = xmap.offset / xmap.elemSize();
int ymap_step = ymap.step / ymap.elemSize(), ymap_offset = ymap.offset / ymap.elemSize();
oclMat coeffsOclMat(coeffsMat.reshape(1, 1));
Context *clCxt = Context::getContext();
string kernelName = "buildWarpAffineMaps";
vector< pair<size_t, const void *> > args;
args.push_back( make_pair( sizeof(cl_mem), (void *)&xmap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&ymap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&coeffsOclMat.data));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap.step));
size_t globalThreads[3] = {xmap.cols, xmap.rows, 1};
size_t localThreads[3] = {32, 8, 1};
openCLExecuteKernel(clCxt, &build_warps, kernelName, globalThreads, localThreads, args, -1, -1);
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_offset));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_offset));
size_t globalThreads[3] = { xmap.cols, xmap.rows, 1 };
size_t localThreads[3] = { 32, 8, 1 };
openCLExecuteKernel(Context::getContext(), &build_warps, "buildWarpAffineMaps", globalThreads, localThreads, args, -1, -1);
}
//////////////////////////////////////////////////////////////////////////////
// buildWarpPerspectiveMaps
void cv::ocl::buildWarpPerspectiveMaps(const Mat &M, bool inverse, Size dsize, oclMat &xmap, oclMat &ymap)
{
CV_Assert(M.rows == 3 && M.cols == 3);
CV_Assert(dsize.area() > 0);
xmap.create(dsize, CV_32FC1);
ymap.create(dsize, CV_32FC1);
@ -235,19 +251,21 @@ void cv::ocl::buildWarpPerspectiveMaps(const Mat &M, bool inverse, Size dsize, o
oclMat coeffsOclMat(coeffsMat.reshape(1, 1));
Context *clCxt = Context::getContext();
string kernelName = "buildWarpPerspectiveMaps";
vector< pair<size_t, const void *> > args;
int xmap_step = xmap.step / xmap.elemSize(), xmap_offset = xmap.offset / xmap.elemSize();
int ymap_step = ymap.step / ymap.elemSize(), ymap_offset = ymap.offset / ymap.elemSize();
vector< pair<size_t, const void *> > args;
args.push_back( make_pair( sizeof(cl_mem), (void *)&xmap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&ymap.data));
args.push_back( make_pair( sizeof(cl_mem), (void *)&coeffsOclMat.data));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.cols));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_step));
args.push_back( make_pair( sizeof(cl_int), (void *)&xmap_offset));
args.push_back( make_pair( sizeof(cl_int), (void *)&ymap_offset));
size_t globalThreads[3] = { xmap.cols, xmap.rows, 1 };
size_t globalThreads[3] = {xmap.cols, xmap.rows, 1};
size_t localThreads[3] = {32, 8, 1};
openCLExecuteKernel(clCxt, &build_warps, kernelName, globalThreads, localThreads, args, -1, -1);
openCLExecuteKernel(Context::getContext(), &build_warps, "buildWarpPerspectiveMaps", globalThreads, NULL, args, -1, -1);
}

149
modules/ocl/src/opencl/build_warps.cl
Unescape View File

@ -43,31 +43,25 @@
//
//M*/
__kernel
void buildWarpPlaneMaps
(
__global float * map_x,
__global float * map_y,
__constant float * KRT,
int tl_u,
int tl_v,
int cols,
int rows,
int step_x,
int step_y,
float scale
)
__kernel void buildWarpPlaneMaps(__global float * xmap, __global float * ymap,
__constant float * KRT,
int tl_u, int tl_v,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset,
float scale)
{
int du = get_global_id(0);
int dv = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
__constant float * ck_rinv = KRT;
__constant float * ct = KRT + 9;
if (du < cols && dv < rows)
{
int xmap_index = mad24(dv, xmap_step, xmap_offset + du);
int ymap_index = mad24(dv, ymap_step, ymap_offset + du);
float u = tl_u + du;
float v = tl_v + dv;
float x, y;
@ -83,33 +77,27 @@ __kernel
x /= z;
y /= z;
map_x[dv * step_x + du] = x;
map_y[dv * step_y + du] = y;
xmap[xmap_index] = x;
ymap[ymap_index] = y;
}
}
__kernel
void buildWarpCylindricalMaps
(
__global float * map_x,
__global float * map_y,
__constant float * ck_rinv,
int tl_u,
int tl_v,
int cols,
int rows,
int step_x,
int step_y,
float scale
)
__kernel void buildWarpCylindricalMaps(__global float * xmap, __global float * ymap,
__constant float * ck_rinv,
int tl_u, int tl_v,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset,
float scale)
{
int du = get_global_id(0);
int dv = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (du < cols && dv < rows)
{
int xmap_index = mad24(dv, xmap_step, xmap_offset + du);
int ymap_index = mad24(dv, ymap_step, ymap_offset + du);
float u = tl_u + du;
float v = tl_v + dv;
float x, y;
@ -127,33 +115,27 @@ __kernel
if (z > 0) { x /= z; y /= z; }
else x = y = -1;
map_x[dv * step_x + du] = x;
map_y[dv * step_y + du] = y;
xmap[xmap_index] = x;
ymap[ymap_index] = y;
}
}
__kernel
void buildWarpSphericalMaps
(
__global float * map_x,
__global float * map_y,
__constant float * ck_rinv,
int tl_u,
int tl_v,
int cols,
int rows,
int step_x,
int step_y,
float scale
)
__kernel void buildWarpSphericalMaps(__global float * xmap, __global float * ymap,
__constant float * ck_rinv,
int tl_u, int tl_v,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset,
float scale)
{
int du = get_global_id(0);
int dv = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (du < cols && dv < rows)
{
int xmap_index = mad24(dv, xmap_step, xmap_offset + du);
int ymap_index = mad24(dv, ymap_step, ymap_offset + du);
float u = tl_u + du;
float v = tl_v + dv;
float x, y;
@ -174,63 +156,52 @@ __kernel
if (z > 0) { x /= z; y /= z; }
else x = y = -1;
map_x[dv * step_x + du] = x;
map_y[dv * step_y + du] = y;
xmap[xmap_index] = x;
ymap[ymap_index] = y;
}
}
__kernel
void buildWarpAffineMaps
(
__global float * xmap,
__global float * ymap,
__constant float * c_warpMat,
int cols,
int rows,
int step_x,
int step_y
)
__kernel void buildWarpAffineMaps(__global float * xmap, __global float * ymap,
__constant float * c_warpMat,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset)
{
int x = get_global_id(0);
int y = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (x < cols && y < rows)
{
const float xcoo = c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2];
const float ycoo = c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5];
int xmap_index = mad24(y, xmap_step, x + xmap_offset);
int ymap_index = mad24(y, ymap_step, x + ymap_offset);
float xcoo = c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2];
float ycoo = c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5];
map_x[y * step_x + x] = xcoo;
map_y[y * step_y + x] = ycoo;
xmap[xmap_index] = xcoo;
ymap[ymap_index] = ycoo;
}
}
__kernel
void buildWarpPerspectiveMaps
(
__global float * xmap,
__global float * ymap,
__constant float * c_warpMat,
int cols,
int rows,
int step_x,
int step_y
)
__kernel void buildWarpPerspectiveMaps(__global float * xmap, __global float * ymap,
__constant float * c_warpMat,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset)
{
int x = get_global_id(0);
int y = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (x < cols && y < rows)
{
const float coeff = 1.0f / (c_warpMat[6] * x + c_warpMat[7] * y + c_warpMat[8]);
int xmap_index = mad24(y, xmap_step, x + xmap_offset);
int ymap_index = mad24(y, ymap_step, x + ymap_offset);
const float xcoo = coeff * (c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2]);
const float ycoo = coeff * (c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5]);
float coeff = 1.0f / (c_warpMat[6] * x + c_warpMat[7] * y + c_warpMat[8]);
float xcoo = coeff * (c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2]);
float ycoo = coeff * (c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5]);
map_x[y * step_x + x] = xcoo;
map_y[y * step_y + x] = ycoo;
xmap[xmap_index] = xcoo;
ymap[ymap_index] = ycoo;
}
}

Write Preview
Loading…
Cancel
Save