// 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);
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);
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);
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);
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);
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);
}
//
//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;
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;
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;
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;
}
}