Merge pull request #2888 from ilya-lavrenov:tapi_remap

modules/imgproc/src/imgwarp.cpp

@@ -3582,7 +3582,9 @@ private:
 static bool ocl_remap(InputArray _src, OutputArray _dst, InputArray _map1, InputArray _map2,
                       int interpolation, int borderType, const Scalar& borderValue)
 {
-    int cn = _src.channels(), type = _src.type(), depth = _src.depth();
+    const ocl::Device & dev = ocl::Device::getDefault();
+    int cn = _src.channels(), type = _src.type(), depth = _src.depth(),
+            rowsPerWI = dev.isIntel() ? 4 : 1;
 
     if (borderType == BORDER_TRANSPARENT || !(interpolation == INTER_LINEAR || interpolation == INTER_NEAREST)
             || _map1.type() == CV_16SC1 || _map2.type() == CV_16SC1)
@@ -3619,12 +3621,14 @@ static bool ocl_remap(InputArray _src, OutputArray _dst, InputArray _map1, Input
     static const char * const interMap[] = { "INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_LINEAR", "INTER_LANCZOS" };
     static const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP",
                            "BORDER_REFLECT_101", "BORDER_TRANSPARENT" };
-    String buildOptions = format("-D %s -D %s -D T=%s", interMap[interpolation], borderMap[borderType], ocl::typeToStr(type));
+    String buildOptions = format("-D %s -D %s -D T=%s -D rowsPerWI=%d",
+                                 interMap[interpolation], borderMap[borderType],
+                                 ocl::typeToStr(type), rowsPerWI);
 
     if (interpolation != INTER_NEAREST)
     {
         char cvt[3][40];
-        int wdepth = std::max(CV_32F, dst.depth());
+        int wdepth = std::max(CV_32F, depth);
         buildOptions = buildOptions
                       + format(" -D WT=%s -D convertToT=%s -D convertToWT=%s"
                                " -D convertToWT2=%s -D WT2=%s",
@@ -3636,10 +3640,9 @@ static bool ocl_remap(InputArray _src, OutputArray _dst, InputArray _map1, Input
     }
     int scalarcn = cn == 3 ? 4 : cn;
     int sctype = CV_MAKETYPE(depth, scalarcn);
-    buildOptions += format(" -D T=%s -D T1=%s"
-                           " -D cn=%d -D ST=%s",
+    buildOptions += format(" -D T=%s -D T1=%s -D cn=%d -D ST=%s -D depth=%d",
                            ocl::typeToStr(type), ocl::typeToStr(depth),
-                           cn, ocl::typeToStr(sctype));
+                           cn, ocl::typeToStr(sctype), depth);
 
     ocl::Kernel k(kernelName.c_str(), ocl::imgproc::remap_oclsrc, buildOptions);
 
@@ -3653,7 +3656,7 @@ static bool ocl_remap(InputArray _src, OutputArray _dst, InputArray _map1, Input
     else
         k.args(srcarg, dstarg, map1arg, ocl::KernelArg::ReadOnlyNoSize(map2), scalararg);
 
-    size_t globalThreads[2] = { dst.cols, dst.rows };
+    size_t globalThreads[2] = { dst.cols, (dst.rows + rowsPerWI - 1) / rowsPerWI };
     return k.run(2, globalThreads, NULL, false);
 }
 

modules/imgproc/src/opencl/remap.cl

@@ -147,16 +147,21 @@ __kernel void remap_2_32FC1(__global const uchar * srcptr, int src_step, int src
                             ST nVal)
 {
     int x = get_global_id(0);
-    int y = get_global_id(1);
+    int y = get_global_id(1) * rowsPerWI;
 
+    if (x < dst_cols)
+    {
         T scalar = convertScalar(nVal);
 
-    if (x < dst_cols && y < dst_rows)
-    {
-        int map1_index = mad24(y, map1_step, x * (int)sizeof(float) + map1_offset);
-        int map2_index = mad24(y, map2_step, x * (int)sizeof(float) + map2_offset);
-        int dst_index = mad24(y, dst_step, x * TSIZE + dst_offset);
+        int map1_index = mad24(y, map1_step, mad24(x, (int)sizeof(float), map1_offset));
+        int map2_index = mad24(y, map2_step, mad24(x, (int)sizeof(float), map2_offset));
+        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
 
+        #pragma unroll
+        for (int i = 0; i < rowsPerWI; ++i, ++y,
+            map1_index += map1_step, map2_index += map2_step, dst_index += dst_step)
+            if (y < dst_rows)
+            {
                 __global const float * map1 = (__global const float *)(map1ptr + map1_index);
                 __global const float * map2 = (__global const float *)(map2ptr + map2_index);
                 __global T * dst = (__global T *)(dstptr + dst_index);
@@ -175,10 +180,11 @@ __kernel void remap_2_32FC1(__global const uchar * srcptr, int src_step, int src
                 }
                 else
                 {
-            int src_index = mad24(gy, src_step, gx * TSIZE + src_offset);
+                    int src_index = mad24(gy, src_step, mad24(gx, TSIZE, src_offset));
                     storepix(loadpix((__global const T*)(srcptr + src_index)), dst);
                 }
             }
+    }
 }
 
 __kernel void remap_32FC2(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
@@ -187,15 +193,19 @@ __kernel void remap_32FC2(__global const uchar * srcptr, int src_step, int src_o
                           ST nVal)
 {
     int x = get_global_id(0);
-    int y = get_global_id(1);
+    int y = get_global_id(1) * rowsPerWI;
 
+    if (x < dst_cols)
+    {
         T scalar = convertScalar(nVal);
+        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
+        int map_index = mad24(y, map_step, mad24(x, (int)sizeof(float2), map_offset));
 
-    if (x < dst_cols && y < dst_rows)
+        #pragma unroll
+        for (int i = 0; i < rowsPerWI; ++i, ++y,
+            map_index += map_step, dst_index += dst_step)
+            if (y < dst_rows)
             {
-        int dst_index = mad24(y, dst_step, x * TSIZE + dst_offset);
-        int map_index = mad24(y, map_step, x * (int)sizeof(float2) + map_offset);
-
                 __global const float2 * map = (__global const float2 *)(mapptr + map_index);
                 __global T * dst = (__global T *)(dstptr + dst_index);
 
@@ -210,10 +220,11 @@ __kernel void remap_32FC2(__global const uchar * srcptr, int src_step, int src_o
                 }
                 else
                 {
-            int src_index = mad24(gy, src_step, gx * TSIZE + src_offset);
+                    int src_index = mad24(gy, src_step, mad24(gx, TSIZE, src_offset));
                     storepix(loadpix((__global const T *)(srcptr + src_index)), dst);
                 }
         }
+    }
 }
 
 __kernel void remap_16SC2(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
@@ -222,15 +233,19 @@ __kernel void remap_16SC2(__global const uchar * srcptr, int src_step, int src_o
                           ST nVal)
 {
     int x = get_global_id(0);
-    int y = get_global_id(1);
+    int y = get_global_id(1) * rowsPerWI;
 
+    if (x < dst_cols)
+    {
         T scalar = convertScalar(nVal);
+        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
+        int map_index = mad24(y, map_step, mad24(x, (int)sizeof(short2), map_offset));
 
-    if (x < dst_cols && y < dst_rows)
+        #pragma unroll
+        for (int i = 0; i < rowsPerWI; ++i, ++y,
+            map_index += map_step, dst_index += dst_step)
+            if (y < dst_rows)
             {
-        int dst_index = mad24(y, dst_step, x * TSIZE + dst_offset);
-        int map_index = mad24(y, map_step, x * (int)sizeof(short2) + map_offset);
-
                 __global const short2 * map = (__global const short2 *)(mapptr + map_index);
                 __global T * dst = (__global T *)(dstptr + dst_index);
 
@@ -245,10 +260,11 @@ __kernel void remap_16SC2(__global const uchar * srcptr, int src_step, int src_o
                 }
                 else
                 {
-            int src_index = mad24(gy, src_step, gx * TSIZE + src_offset);
+                    int src_index = mad24(gy, src_step, mad24(gx, TSIZE, src_offset));
                     storepix(loadpix((__global const T *)(srcptr + src_index)), dst);
                 }
             }
+    }
 }
 
 __kernel void remap_16SC2_16UC1(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
@@ -258,16 +274,20 @@ __kernel void remap_16SC2_16UC1(__global const uchar * srcptr, int src_step, int
                                 ST nVal)
 {
     int x = get_global_id(0);
-    int y = get_global_id(1);
+    int y = get_global_id(1) * rowsPerWI;
 
+    if (x < dst_cols)
+    {
         T scalar = convertScalar(nVal);
-
-    if (x < dst_cols && y < dst_rows)
+        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
+        int map1_index = mad24(y, map1_step, mad24(x, (int)sizeof(short2), map1_offset));
+        int map2_index = mad24(y, map2_step, mad24(x, (int)sizeof(ushort), map2_offset));
+
+        #pragma unroll
+        for (int i = 0; i < rowsPerWI; ++i, ++y,
+            map1_index += map1_step, map2_index += map2_step, dst_index += dst_step)
+            if (y < dst_rows)
             {
-        int dst_index = mad24(y, dst_step, x * TSIZE + dst_offset);
-        int map1_index = mad24(y, map1_step, x * (int)sizeof(short2) + map1_offset);
-        int map2_index = mad24(y, map2_step, x * (int)sizeof(ushort) + map2_offset);
-
                 __global const short2 * map1 = (__global const short2 *)(map1ptr + map1_index);
                 __global const ushort * map2 = (__global const ushort *)(map2ptr + map2_index);
                 __global T * dst = (__global T *)(dstptr + dst_index);
@@ -286,13 +306,22 @@ __kernel void remap_16SC2_16UC1(__global const uchar * srcptr, int src_step, int
                 }
                 else
                 {
-            int src_index = mad24(gy, src_step, gx * TSIZE + src_offset);
+                    int src_index = mad24(gy, src_step, mad24(gx, TSIZE, src_offset));
                     storepix(loadpix((__global const T *)(srcptr + src_index)), dst);
                 }
             }
+    }
 }
 
-#elif INTER_LINEAR
+#elif defined INTER_LINEAR
+
+__constant float coeffs[64] =
+{ 1.000000f, 0.000000f, 0.968750f, 0.031250f, 0.937500f, 0.062500f, 0.906250f, 0.093750f, 0.875000f, 0.125000f, 0.843750f, 0.156250f,
+  0.812500f, 0.187500f, 0.781250f, 0.218750f, 0.750000f, 0.250000f, 0.718750f, 0.281250f, 0.687500f, 0.312500f, 0.656250f, 0.343750f,
+  0.625000f, 0.375000f, 0.593750f, 0.406250f, 0.562500f, 0.437500f, 0.531250f, 0.468750f, 0.500000f, 0.500000f, 0.468750f, 0.531250f,
+  0.437500f, 0.562500f, 0.406250f, 0.593750f, 0.375000f, 0.625000f, 0.343750f, 0.656250f, 0.312500f, 0.687500f, 0.281250f, 0.718750f,
+  0.250000f, 0.750000f, 0.218750f, 0.781250f, 0.187500f, 0.812500f, 0.156250f, 0.843750f, 0.125000f, 0.875000f, 0.093750f, 0.906250f,
+  0.062500f, 0.937500f, 0.031250f, 0.968750f };
 
 __kernel void remap_16SC2_16UC1(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
                                 __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
@@ -301,14 +330,20 @@ __kernel void remap_16SC2_16UC1(__global const uchar * srcptr, int src_step, int
                                 ST nVal)
 {
     int x = get_global_id(0);
-    int y = get_global_id(1);
+    int y = get_global_id(1) * rowsPerWI;
 
-    if (x < dst_cols && y < dst_rows)
+    if (x < dst_cols)
+    {
+        WT scalar = convertToWT(convertScalar(nVal));
+        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
+        int map1_index = mad24(y, map1_step, mad24(x, (int)sizeof(short2), map1_offset));
+        int map2_index = mad24(y, map2_step, mad24(x, (int)sizeof(ushort), map2_offset));
+
+        #pragma unroll
+        for (int i = 0; i < rowsPerWI; ++i, ++y,
+            map1_index += map1_step, map2_index += map2_step, dst_index += dst_step)
+            if (y < dst_rows)
             {
-        int dst_index = mad24(y, dst_step, x * TSIZE + dst_offset);
-        int map1_index = mad24(y, map1_step, x * (int)sizeof(short2) + map1_offset);
-        int map2_index = mad24(y, map2_step, x * (int)sizeof(ushort) + map2_offset);
-
                 __global const short2 * map1 = (__global const short2 *)(map1ptr + map1_index);
                 __global const ushort * map2 = (__global const ushort *)(map2ptr + map2_index);
                 __global T * dst = (__global T *)(dstptr + dst_index);
@@ -321,7 +356,6 @@ __kernel void remap_16SC2_16UC1(__global const uchar * srcptr, int src_step, int
                 ushort map2Value = (ushort)(map2[0] & (INTER_TAB_SIZE2 - 1));
                 WT2 u = (WT2)(map2Value & (INTER_TAB_SIZE - 1), map2Value >> INTER_BITS) / (WT2)(INTER_TAB_SIZE);
 
-        WT scalar = convertToWT(convertScalar(nVal));
                 WT a = scalar, b = scalar, c = scalar, d = scalar;
 
                 if (!NEED_EXTRAPOLATION(map_dataA.x, map_dataA.y))
@@ -350,6 +384,7 @@ __kernel void remap_16SC2_16UC1(__global const uchar * srcptr, int src_step, int
                               d * (u.x)     * (u.y);
                 storepix(convertToT(dst_data), dst);
             }
+    }
 }
 
 __kernel void remap_2_32FC1(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
@@ -359,18 +394,67 @@ __kernel void remap_2_32FC1(__global const uchar * srcptr, int src_step, int src
                             ST nVal)
 {
     int x = get_global_id(0);
-    int y = get_global_id(1);
+    int y = get_global_id(1) * rowsPerWI;
 
-    if (x < dst_cols && y < dst_rows)
+    if (x < dst_cols)
+    {
+        WT scalar = convertToWT(convertScalar(nVal));
+        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
+        int map1_index = mad24(y, map1_step, mad24(x, (int)sizeof(float), map1_offset));
+        int map2_index = mad24(y, map2_step, mad24(x, (int)sizeof(float), map2_offset));
+
+        #pragma unroll
+        for (int i = 0; i < rowsPerWI; ++i, ++y,
+            map1_index += map1_step, map2_index += map2_step, dst_index += dst_step)
+            if (y < dst_rows)
             {
-        int dst_index = mad24(y, dst_step, x * TSIZE + dst_offset);
-        int map1_index = mad24(y, map1_step, x * (int)sizeof(float) + map1_offset);
-        int map2_index = mad24(y, map2_step, x * (int)sizeof(float) + map2_offset);
-
                 __global const float * map1 = (__global const float *)(map1ptr + map1_index);
                 __global const float * map2 = (__global const float *)(map2ptr + map2_index);
                 __global T * dst = (__global T *)(dstptr + dst_index);
 
+#if defined BORDER_CONSTANT
+
+                float xf = map1[0], yf = map2[0];
+                int sx = convert_int_sat_rtn(xf), sy = convert_int_sat_rtn(yf);
+
+                __constant float * coeffs_x = coeffs + ((convert_int_rte(xf * INTER_TAB_SIZE) & (INTER_TAB_SIZE - 1)) << 1);
+                __constant float * coeffs_y = coeffs + ((convert_int_rte(yf * INTER_TAB_SIZE) & (INTER_TAB_SIZE - 1)) << 1);
+
+                WT sum = (WT)(0), xsum;
+                int src_index = mad24(sy, src_step, mad24(sx, TSIZE, src_offset));
+
+                #pragma unroll
+                for (int yp = 0; yp < 2; ++yp, src_index += src_step)
+                {
+                    if (sy + yp >= 0 && sy + yp < src_rows)
+                    {
+                        xsum = (WT)(0);
+                        if (sx >= 0 && sx + 2 < src_cols)
+                        {
+#if depth == 0 && cn == 1
+                            uchar2 value = vload2(0, srcptr + src_index);
+                            xsum = dot(convert_float2(value), (float2)(coeffs_x[0], coeffs_x[1]));
+#else
+                            #pragma unroll
+                            for (int xp = 0; xp < 2; ++xp)
+                                xsum = fma(convertToWT(loadpix(srcptr + mad24(xp, TSIZE, src_index))), coeffs_x[xp], xsum);
+#endif
+                        }
+                        else
+                        {
+                            #pragma unroll
+                            for (int xp = 0; xp < 2; ++xp)
+                                xsum = fma(sx + xp >= 0 && sx + xp < src_cols ?
+                                           convertToWT(loadpix(srcptr + mad24(xp, TSIZE, src_index))) : scalar, coeffs_x[xp], xsum);
+                        }
+                        sum = fma(xsum, coeffs_y[yp], sum);
+                    }
+                    else
+                        sum = fma(scalar, coeffs_y[yp], sum);
+                }
+
+                storepix(convertToT(sum), dst);
+#else
                 float2 map_data = (float2)(map1[0], map2[0]);
 
                 int2 map_dataA = convert_int2_sat_rtn(map_data);
@@ -408,6 +492,8 @@ __kernel void remap_2_32FC1(__global const uchar * srcptr, int src_step, int src
                               c * (1 - u.x) * (u.y) +
                               d * (u.x)     * (u.y);
                 storepix(convertToT(dst_data), dst);
+#endif
+            }
     }
 }
 
@@ -417,13 +503,19 @@ __kernel void remap_32FC2(__global const uchar * srcptr, int src_step, int src_o
                           ST nVal)
 {
     int x = get_global_id(0);
-    int y = get_global_id(1);
+    int y = get_global_id(1) * rowsPerWI;
 
-    if (x < dst_cols && y < dst_rows)
+    if (x < dst_cols)
     {
-        int dst_index = mad24(y, dst_step, x * TSIZE + dst_offset);
-        int map_index = mad24(y, map_step, x * (int)sizeof(float2) + map_offset);
+        WT scalar = convertToWT(convertScalar(nVal));
+        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
+        int map_index = mad24(y, map_step, mad24(x, (int)sizeof(float2), map_offset));
 
+        #pragma unroll
+        for (int i = 0; i < rowsPerWI; ++i, ++y,
+            map_index += map_step, dst_index += dst_step)
+            if (y < dst_rows)
+            {
                 __global const float2 * map = (__global const float2 *)(mapptr + map_index);
                 __global T * dst = (__global T *)(dstptr + dst_index);
 
@@ -435,7 +527,6 @@ __kernel void remap_32FC2(__global const uchar * srcptr, int src_step, int src_o
 
                 float2 _u = map_data - convert_float2(map_dataA);
                 WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)INTER_TAB_SIZE)) / (WT2)INTER_TAB_SIZE;
-        WT scalar = convertToWT(convertScalar(nVal));
                 WT a = scalar, b = scalar, c = scalar, d = scalar;
 
                 if (!NEED_EXTRAPOLATION(map_dataA.x, map_dataA.y))
@@ -464,6 +555,7 @@ __kernel void remap_32FC2(__global const uchar * srcptr, int src_step, int src_o
                               d * (u.x)     * (u.y);
                 storepix(convertToT(dst_data), dst);
             }
+    }
 }
 
 #endif

modules/imgproc/test/ocl/test_warp.cpp

@@ -267,7 +267,7 @@ PARAM_TEST_CASE(Remap, MatDepth, Channels, std::pair<MatType, MatType>, BorderTy
         Border map1Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
         randomSubMat(map1, map1_roi, dstROISize, map1Border, map1Type, -mapMaxValue, mapMaxValue);
 
-        Border map2Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        Border map2Border = randomBorder(0, useRoi ? MAX_VALUE + 1 : 0);
         if (map2Type != noType)
         {
             int mapMinValue = -mapMaxValue;