added ROI support to ocl::buildWarp*Maps functions

diff --git a/modules/ocl/src/build_warps.cpp b/modules/ocl/src/build_warps.cpp
index dc9ab66..40c082b 100644 (file)
--- 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<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);
 }

diff --git a/modules/ocl/src/opencl/build_warps.cl b/modules/ocl/src/opencl/build_warps.cl
index 4402e8c..bd5e002 100644 (file)
--- 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);
+
+        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;
     }
 }