From 62d4393883b21b55b465ace2df232405c66c22f2 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Wed, 22 Oct 2014 17:43:29 +0300 Subject: [PATCH 1/3] Small optimization for buildMaps kernels in case rowsPerWI > 1 --- modules/stitching/src/opencl/warpers.cl | 39 +++++++++++-------------- modules/stitching/src/warpers.cpp | 6 ++-- 2 files changed, 20 insertions(+), 25 deletions(-) diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl index 9b5619fcad..7ec87ae2f8 100644 --- a/modules/stitching/src/opencl/warpers.cl +++ b/modules/stitching/src/opencl/warpers.cl @@ -56,19 +56,19 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset)); int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); + float u = tl_u + du; + float x_ = u * scale - ct[0]; + float ct1 = 1 - ct[2]; + for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, ymap_index += ymap_step) { __global float * xmap = (__global float *)(xmapptr + xmap_index); __global float * ymap = (__global float *)(ymapptr + ymap_index); - float u = tl_u + du; float v = tl_v + dv; + float y_ = v * scale - ct[1]; - float x_ = u / scale - ct[0]; - float y_ = v / scale - ct[1]; - - float ct1 = 1 - ct[2]; float x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * ct1)); float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1)); float z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * ct1)); @@ -94,22 +94,19 @@ __kernel void buildWarpCylindricalMaps(__global uchar * xmapptr, int xmap_step, int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset)); int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); + float u = (tl_u + du) * scale; + float x_, z_; + x_ = sincos(u, &z_); + for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, ymap_index += ymap_step) { __global float * xmap = (__global float *)(xmapptr + xmap_index); __global float * ymap = (__global float *)(ymapptr + ymap_index); - float u = tl_u + du; - float v = tl_v + dv; - float x, y; + float y_ = (tl_v + dv) * scale; - u /= scale; - float x_, y_, z_; - x_ = sincos(u, &z_); - y_ = v / scale; - - float z; + float x, y, z; x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_)); y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_)); z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_)); @@ -137,25 +134,23 @@ __kernel void buildWarpSphericalMaps(__global uchar * xmapptr, int xmap_step, in int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset)); int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); + float u = (tl_u + du) * scale; + float cosu, sinu = sincos(u, &cosu); + for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, ymap_index += ymap_step) { __global float * xmap = (__global float *)(xmapptr + xmap_index); __global float * ymap = (__global float *)(ymapptr + ymap_index); - float u = tl_u + du; - float v = tl_v + dv; - float x, y; + float v = (tl_v + dv) * scale; - v /= scale; - u /= scale; - - float cosv, sinv = sincos(v, &cosv), cosu, sinu = sincos(u, &cosu); + float cosv, sinv = sincos(v, &cosv); float x_ = sinv * sinu; float y_ = -cosv; float z_ = sinv * cosu; - float z; + float x, y, z; x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_)); y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_)); z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_)); diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index f474010ac5..2711a012f6 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -110,7 +110,7 @@ Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArra k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), ocl::KernelArg::PtrReadOnly(uk_rinv), ocl::KernelArg::PtrReadOnly(ut), - dst_tl.x, dst_tl.y, projector_.scale, rowsPerWI); + dst_tl.x, dst_tl.y, 1/projector_.scale, rowsPerWI); size_t globalsize[2] = { dsize.width, (dsize.height + rowsPerWI - 1) / rowsPerWI }; if (k.run(2, globalsize, NULL, true)) @@ -388,7 +388,7 @@ Rect SphericalWarper::buildMaps(Size src_size, InputArray K, InputArray R, Outpu UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), - ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale, rowsPerWI); + ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, 1/projector_.scale, rowsPerWI); size_t globalsize[2] = { dsize.width, (dsize.height + rowsPerWI - 1) / rowsPerWI }; if (k.run(2, globalsize, NULL, true)) @@ -436,7 +436,7 @@ Rect CylindricalWarper::buildMaps(Size src_size, InputArray K, InputArray R, Out UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), - ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale, + ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, 1/projector_.scale, rowsPerWI); size_t globalsize[2] = { dsize.width, (dsize.height + rowsPerWI - 1) / rowsPerWI }; From 0aab7795326c5a56039c6a7eed7d4bb7f78277e3 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Fri, 24 Oct 2014 11:03:10 +0300 Subject: [PATCH 2/3] Overload PlaneWarper::buildMaps method from base class --- .../stitching/include/opencv2/stitching/detail/warpers.hpp | 1 + modules/stitching/src/opencl/warpers.cl | 4 ++-- modules/stitching/src/warpers.cpp | 5 +++++ 3 files changed, 8 insertions(+), 2 deletions(-) diff --git a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp index c8869f116d..ac9e256ad7 100644 --- a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp @@ -135,6 +135,7 @@ public: Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R, InputArray T); virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap); virtual Point warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, OutputArray dst); diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl index 7ec87ae2f8..9d245893cf 100644 --- a/modules/stitching/src/opencl/warpers.cl +++ b/modules/stitching/src/opencl/warpers.cl @@ -57,7 +57,7 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset)); float u = tl_u + du; - float x_ = u * scale - ct[0]; + float x_ = fma(u, scale, -ct[0]); float ct1 = 1 - ct[2]; for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step, @@ -67,7 +67,7 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm __global float * ymap = (__global float *)(ymapptr + ymap_index); float v = tl_v + dv; - float y_ = v * scale - ct[1]; + float y_ = fma(v, scale, -ct[1]); float x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * ct1)); float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1)); diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index 2711a012f6..744474ba6e 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -87,6 +87,11 @@ Point2f PlaneWarper::warpPoint(const Point2f &pt, InputArray K, InputArray R, In return uv; } +Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) +{ + return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32FC1), xmap, ymap); +} + Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray _xmap, OutputArray _ymap) { projector_.setCameraParams(K, R, T); From ebfaf4c5d871397d7de10f791907b2e7e0246fbd Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Mon, 27 Oct 2014 15:38:44 +0300 Subject: [PATCH 3/3] Added checking that z is non zero to buildWarpPlaneMaps kernel --- modules/stitching/src/opencl/warpers.cl | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl index 9d245893cf..a8f04a63cd 100644 --- a/modules/stitching/src/opencl/warpers.cl +++ b/modules/stitching/src/opencl/warpers.cl @@ -73,8 +73,10 @@ __kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xm float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1)); float z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * ct1)); - x /= z; - y /= z; + if (z != 0) + x /= z, y /= z; + else + x = y = -1; xmap[0] = x; ymap[0] = y;