From 4248f8221146a929b1ac98f24d58c699e8dafa03 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 8 Nov 2013 18:40:53 +0400 Subject: [PATCH] added ROI support to ocl::buildWarp*Maps functions --- modules/ocl/src/build_warps.cpp | 144 ++++++++++++++----------- modules/ocl/src/opencl/build_warps.cl | 149 +++++++++++--------------- 2 files changed, 141 insertions(+), 152 deletions(-) diff --git a/modules/ocl/src/build_warps.cpp b/modules/ocl/src/build_warps.cpp index dc9ab66db9..40c082b556 100644 --- a/modules/ocl/src/build_warps.cpp +++ b/modules/ocl/src/build_warps.cpp @@ -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 > 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 > 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 > 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 > 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 > 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 > 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 > 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 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, "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 > 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 > 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 localThreads[3] = {32, 8, 1}; - openCLExecuteKernel(clCxt, &build_warps, kernelName, globalThreads, localThreads, args, -1, -1); + size_t globalThreads[3] = { xmap.cols, xmap.rows, 1 }; + + openCLExecuteKernel(Context::getContext(), &build_warps, "buildWarpPerspectiveMaps", globalThreads, NULL, args, -1, -1); } diff --git a/modules/ocl/src/opencl/build_warps.cl b/modules/ocl/src/opencl/build_warps.cl index 4402e8c388..bd5e002b5d 100644 --- a/modules/ocl/src/opencl/build_warps.cl +++ b/modules/ocl/src/opencl/build_warps.cl @@ -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); - map_x[y * step_x + x] = xcoo; - map_y[y * step_y + x] = ycoo; + 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]; + + 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; } }