From 60b59eeb73ebe1c719181db1aa4aaa43e1df2c1a Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 20 Nov 2013 15:40:56 +0400 Subject: [PATCH] workaround for AMD bug: UNREACHABLE EXECUTED --- modules/ocl/src/arithm.cpp | 7 +++++++ modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl | 6 ++++++ 2 files changed, 13 insertions(+) diff --git a/modules/ocl/src/arithm.cpp b/modules/ocl/src/arithm.cpp index 8a2390d..9bd09d6 100644 --- a/modules/ocl/src/arithm.cpp +++ b/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(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 )); diff --git a/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl b/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl index e03fa69..e07f314 100644 --- a/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl +++ b/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 -- 2.7.4