Merge pull request #1836 from ilya-lavrenov:ocl_norm_workaround

modules/ocl/src/arithm.cpp

@@ -603,6 +603,12 @@ static void arithm_absdiff_nonsaturate_run(const oclMat & src1, const oclMat & s
     }
     CV_Assert(src1.step % src1.elemSize() == 0 && (src2.empty() || src2.step % src2.elemSize() == 0));
 
+    if (src2.empty() && (src1.depth() == CV_8U || src1.depth() == CV_16U))
+    {
+        src1.convertTo(diff, CV_32S);
+        return;
+    }
+
     int ddepth = std::max(src1.depth(), CV_32S);
     if (ntype == NORM_L2)
         ddepth = std::max<int>(CV_32F, ddepth);
@@ -639,6 +645,7 @@ static void arithm_absdiff_nonsaturate_run(const oclMat & src1, const oclMat & s
         args.push_back( make_pair( sizeof(cl_int), (void *)&src2offset1 ));
 
         kernelName += "_binary";
+        buildOptions += " -D BINARY";
     }
 
     args.push_back( make_pair( sizeof(cl_mem), (void *)&diff.data ));

modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl

@@ -52,6 +52,8 @@
 #endif
 #endif
 
+#ifdef BINARY
+
 __kernel void arithm_absdiff_nonsaturate_binary(__global srcT *src1, int src1_step, int src1_offset,
                          __global srcT *src2, int src2_step, int src2_offset,
                          __global dstT *dst, int dst_step, int dst_offset,
@@ -78,6 +80,8 @@ __kernel void arithm_absdiff_nonsaturate_binary(__global srcT *src1, int src1_st
     }
 }
 
+#else
+
 __kernel void arithm_absdiff_nonsaturate(__global srcT *src1, int src1_step, int src1_offset,
                          __global dstT *dst, int dst_step, int dst_offset,
                          int cols, int rows)
@@ -99,3 +103,5 @@ __kernel void arithm_absdiff_nonsaturate(__global srcT *src1, int src1_step, int
         }
     }
 }
+
+#endif