From a5f68e98f4ec18036ef800482127f0f72b0e26a7 Mon Sep 17 00:00:00 2001 From: k-shinotsuka Date: Mon, 30 Apr 2018 16:38:49 +0900 Subject: [PATCH 01/14] add universal intrinsics for HLS2RGB_f --- modules/imgproc/src/color_hsv.cpp | 198 +++++++++++------------------- 1 file changed, 71 insertions(+), 127 deletions(-) diff --git a/modules/imgproc/src/color_hsv.cpp b/modules/imgproc/src/color_hsv.cpp index 45f5bdb21d..81530289dd 100644 --- a/modules/imgproc/src/color_hsv.cpp +++ b/modules/imgproc/src/color_hsv.cpp @@ -872,167 +872,111 @@ struct HLS2RGB_f HLS2RGB_f(int _dstcn, int _blueIdx, float _hrange) : dstcn(_dstcn), blueIdx(_blueIdx), hscale(6.f/_hrange) { - #if CV_SSE2 - haveSIMD = checkHardwareSupport(CV_CPU_SSE2); + #if CV_SIMD128 + hasSIMD = hasSIMD128(); #endif } - #if CV_SSE2 - void process(__m128& v_h0, __m128& v_h1, __m128& v_l0, - __m128& v_l1, __m128& v_s0, __m128& v_s1) const + #if CV_SIMD128 + inline void process(v_float32x4& v_h, v_float32x4& v_l, v_float32x4& v_s) const { - __m128 v_lel0 = _mm_cmple_ps(v_l0, _mm_set1_ps(0.5f)); - __m128 v_lel1 = _mm_cmple_ps(v_l1, _mm_set1_ps(0.5f)); - __m128 v_p20 = _mm_andnot_ps(v_lel0, _mm_sub_ps(_mm_add_ps(v_l0, v_s0), _mm_mul_ps(v_l0, v_s0))); - __m128 v_p21 = _mm_andnot_ps(v_lel1, _mm_sub_ps(_mm_add_ps(v_l1, v_s1), _mm_mul_ps(v_l1, v_s1))); - v_p20 = _mm_or_ps(v_p20, _mm_and_ps(v_lel0, _mm_mul_ps(v_l0, _mm_add_ps(_mm_set1_ps(1.0f), v_s0)))); - v_p21 = _mm_or_ps(v_p21, _mm_and_ps(v_lel1, _mm_mul_ps(v_l1, _mm_add_ps(_mm_set1_ps(1.0f), v_s1)))); + v_float32x4 v_one = v_setall_f32(1.0f); - __m128 v_p10 = _mm_sub_ps(_mm_mul_ps(_mm_set1_ps(2.0f), v_l0), v_p20); - __m128 v_p11 = _mm_sub_ps(_mm_mul_ps(_mm_set1_ps(2.0f), v_l1), v_p21); + v_float32x4 v_l_le_half = v_l <= v_setall_f32(0.5f); + v_float32x4 v_ls = v_l * v_s; + v_float32x4 v_elem0 = v_select(v_l_le_half, v_ls, v_s - v_ls); - v_h0 = _mm_mul_ps(v_h0, _mm_set1_ps(hscale)); - v_h1 = _mm_mul_ps(v_h1, _mm_set1_ps(hscale)); + v_float32x4 v_hs_raw = v_h * v_setall_f32(hscale); + v_float32x4 v_pre_hs = v_cvt_f32(v_trunc(v_hs_raw)); + v_float32x4 v_hs = v_hs_raw - v_pre_hs; + v_float32x4 v_sector = v_pre_hs - v_setall_f32(6.0f) * v_cvt_f32(v_trunc(v_hs_raw * v_setall_f32(1.0f / 6.0f))); + v_float32x4 v_elem1 = v_hs + v_hs; - __m128 v_pre_sector0 = _mm_cvtepi32_ps(_mm_cvttps_epi32(v_h0)); - __m128 v_pre_sector1 = _mm_cvtepi32_ps(_mm_cvttps_epi32(v_h1)); + v_float32x4 v_tab0 = v_l + v_elem0; + v_float32x4 v_tab1 = v_l - v_elem0; + v_float32x4 v_tab2 = v_l + v_elem0 - v_elem0 * v_elem1; + v_float32x4 v_tab3 = v_l - v_elem0 + v_elem0 * v_elem1; - v_h0 = _mm_sub_ps(v_h0, v_pre_sector0); - v_h1 = _mm_sub_ps(v_h1, v_pre_sector1); + v_float32x4 v_two = v_setall_f32(2.0f); + v_float32x4 v_four = v_setall_f32(4.0f); - __m128 v_p2_p10 = _mm_sub_ps(v_p20, v_p10); - __m128 v_p2_p11 = _mm_sub_ps(v_p21, v_p11); - __m128 v_tab20 = _mm_add_ps(v_p10, _mm_mul_ps(v_p2_p10, _mm_sub_ps(_mm_set1_ps(1.0f), v_h0))); - __m128 v_tab21 = _mm_add_ps(v_p11, _mm_mul_ps(v_p2_p11, _mm_sub_ps(_mm_set1_ps(1.0f), v_h1))); - __m128 v_tab30 = _mm_add_ps(v_p10, _mm_mul_ps(v_p2_p10, v_h0)); - __m128 v_tab31 = _mm_add_ps(v_p11, _mm_mul_ps(v_p2_p11, v_h1)); + v_h = v_select(v_sector < v_two , v_tab1, + v_select(v_sector <= v_two , v_tab3, + v_select(v_sector <= v_four, v_tab0, v_tab2))); - __m128 v_sector0 = _mm_div_ps(v_pre_sector0, _mm_set1_ps(6.0f)); - __m128 v_sector1 = _mm_div_ps(v_pre_sector1, _mm_set1_ps(6.0f)); - v_sector0 = _mm_cvtepi32_ps(_mm_cvttps_epi32(v_sector0)); - v_sector1 = _mm_cvtepi32_ps(_mm_cvttps_epi32(v_sector1)); - v_sector0 = _mm_mul_ps(v_sector0, _mm_set1_ps(6.0f)); - v_sector1 = _mm_mul_ps(v_sector1, _mm_set1_ps(6.0f)); - v_sector0 = _mm_sub_ps(v_pre_sector0, v_sector0); - v_sector1 = _mm_sub_ps(v_pre_sector1, v_sector1); + v_l = v_select(v_sector < v_one , v_tab3, + v_select(v_sector <= v_two , v_tab0, + v_select(v_sector < v_four, v_tab2, v_tab1))); - v_h0 = _mm_and_ps(v_p10, _mm_cmplt_ps(v_sector0, _mm_set1_ps(2.0f))); - v_h1 = _mm_and_ps(v_p11, _mm_cmplt_ps(v_sector1, _mm_set1_ps(2.0f))); - v_h0 = _mm_or_ps(v_h0, _mm_and_ps(v_tab30, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(2.0f)))); - v_h1 = _mm_or_ps(v_h1, _mm_and_ps(v_tab31, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(2.0f)))); - v_h0 = _mm_or_ps(v_h0, _mm_and_ps(v_p20, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(3.0f)))); - v_h1 = _mm_or_ps(v_h1, _mm_and_ps(v_p21, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(3.0f)))); - v_h0 = _mm_or_ps(v_h0, _mm_and_ps(v_p20, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(4.0f)))); - v_h1 = _mm_or_ps(v_h1, _mm_and_ps(v_p21, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(4.0f)))); - v_h0 = _mm_or_ps(v_h0, _mm_and_ps(v_tab20, _mm_cmpgt_ps(v_sector0, _mm_set1_ps(4.0f)))); - v_h1 = _mm_or_ps(v_h1, _mm_and_ps(v_tab21, _mm_cmpgt_ps(v_sector1, _mm_set1_ps(4.0f)))); - v_l0 = _mm_and_ps(v_tab30, _mm_cmplt_ps(v_sector0, _mm_set1_ps(1.0f))); - v_l1 = _mm_and_ps(v_tab31, _mm_cmplt_ps(v_sector1, _mm_set1_ps(1.0f))); - v_l0 = _mm_or_ps(v_l0, _mm_and_ps(v_p20, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(1.0f)))); - v_l1 = _mm_or_ps(v_l1, _mm_and_ps(v_p21, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(1.0f)))); - v_l0 = _mm_or_ps(v_l0, _mm_and_ps(v_p20, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(2.0f)))); - v_l1 = _mm_or_ps(v_l1, _mm_and_ps(v_p21, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(2.0f)))); - v_l0 = _mm_or_ps(v_l0, _mm_and_ps(v_tab20, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(3.0f)))); - v_l1 = _mm_or_ps(v_l1, _mm_and_ps(v_tab21, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(3.0f)))); - v_l0 = _mm_or_ps(v_l0, _mm_and_ps(v_p10, _mm_cmpgt_ps(v_sector0, _mm_set1_ps(3.0f)))); - v_l1 = _mm_or_ps(v_l1, _mm_and_ps(v_p11, _mm_cmpgt_ps(v_sector1, _mm_set1_ps(3.0f)))); - v_s0 = _mm_and_ps(v_p20, _mm_cmplt_ps(v_sector0, _mm_set1_ps(1.0f))); - v_s1 = _mm_and_ps(v_p21, _mm_cmplt_ps(v_sector1, _mm_set1_ps(1.0f))); - v_s0 = _mm_or_ps(v_s0, _mm_and_ps(v_tab20, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(1.0f)))); - v_s1 = _mm_or_ps(v_s1, _mm_and_ps(v_tab21, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(1.0f)))); - v_s0 = _mm_or_ps(v_s0, _mm_and_ps(v_p10, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(2.0f)))); - v_s1 = _mm_or_ps(v_s1, _mm_and_ps(v_p11, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(2.0f)))); - v_s0 = _mm_or_ps(v_s0, _mm_and_ps(v_p10, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(3.0f)))); - v_s1 = _mm_or_ps(v_s1, _mm_and_ps(v_p11, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(3.0f)))); - v_s0 = _mm_or_ps(v_s0, _mm_and_ps(v_tab30, _mm_cmpeq_ps(v_sector0, _mm_set1_ps(4.0f)))); - v_s1 = _mm_or_ps(v_s1, _mm_and_ps(v_tab31, _mm_cmpeq_ps(v_sector1, _mm_set1_ps(4.0f)))); - v_s0 = _mm_or_ps(v_s0, _mm_and_ps(v_p20, _mm_cmpgt_ps(v_sector0, _mm_set1_ps(4.0f)))); - v_s1 = _mm_or_ps(v_s1, _mm_and_ps(v_p21, _mm_cmpgt_ps(v_sector1, _mm_set1_ps(4.0f)))); + v_s = v_select(v_sector < v_one , v_tab0, + v_select(v_sector < v_two , v_tab2, + v_select(v_sector < v_four, v_tab1, + v_select(v_sector <= v_four, v_tab3, v_tab0)))); } #endif void operator()(const float* src, float* dst, int n) const { int i = 0, bidx = blueIdx, dcn = dstcn; - float _hscale = hscale; float alpha = ColorChannel::max(); n *= 3; - #if CV_SSE2 - if (haveSIMD) + #if CV_SIMD128 + if (hasSIMD) { - for( ; i <= n - 24; i += 24, dst += dcn * 8 ) + if (dcn == 3) { - __m128 v_h0 = _mm_loadu_ps(src + i + 0); - __m128 v_h1 = _mm_loadu_ps(src + i + 4); - __m128 v_l0 = _mm_loadu_ps(src + i + 8); - __m128 v_l1 = _mm_loadu_ps(src + i + 12); - __m128 v_s0 = _mm_loadu_ps(src + i + 16); - __m128 v_s1 = _mm_loadu_ps(src + i + 20); - - _mm_deinterleave_ps(v_h0, v_h1, v_l0, v_l1, v_s0, v_s1); - - process(v_h0, v_h1, v_l0, v_l1, v_s0, v_s1); - - if (dcn == 3) + if (bidx) { - if (bidx) + for (; i <= n - 12; i += 12, dst += dcn * 4) { - _mm_interleave_ps(v_s0, v_s1, v_l0, v_l1, v_h0, v_h1); - - _mm_storeu_ps(dst + 0, v_s0); - _mm_storeu_ps(dst + 4, v_s1); - _mm_storeu_ps(dst + 8, v_l0); - _mm_storeu_ps(dst + 12, v_l1); - _mm_storeu_ps(dst + 16, v_h0); - _mm_storeu_ps(dst + 20, v_h1); + v_float32x4 v_h; + v_float32x4 v_l; + v_float32x4 v_s; + v_load_deinterleave(src + i, v_h, v_l, v_s); + process(v_h, v_l, v_s); + v_store_interleave(dst, v_s, v_l, v_h); } - else + } else { + for (; i <= n - 12; i += 12, dst += dcn * 4) { - _mm_interleave_ps(v_h0, v_h1, v_l0, v_l1, v_s0, v_s1); - - _mm_storeu_ps(dst + 0, v_h0); - _mm_storeu_ps(dst + 4, v_h1); - _mm_storeu_ps(dst + 8, v_l0); - _mm_storeu_ps(dst + 12, v_l1); - _mm_storeu_ps(dst + 16, v_s0); - _mm_storeu_ps(dst + 20, v_s1); + v_float32x4 v_h; + v_float32x4 v_l; + v_float32x4 v_s; + v_load_deinterleave(src + i, v_h, v_l, v_s); + process(v_h, v_l, v_s); + v_store_interleave(dst, v_h, v_l, v_s); } } - else + } else { // dcn == 4 + if (bidx) { - __m128 v_a0 = _mm_set1_ps(alpha); - __m128 v_a1 = _mm_set1_ps(alpha); - if (bidx) + for (; i <= n - 12; i += 12, dst += dcn * 4) { - _mm_interleave_ps(v_s0, v_s1, v_l0, v_l1, v_h0, v_h1, v_a0, v_a1); - - _mm_storeu_ps(dst + 0, v_s0); - _mm_storeu_ps(dst + 4, v_s1); - _mm_storeu_ps(dst + 8, v_l0); - _mm_storeu_ps(dst + 12, v_l1); - _mm_storeu_ps(dst + 16, v_h0); - _mm_storeu_ps(dst + 20, v_h1); - _mm_storeu_ps(dst + 24, v_a0); - _mm_storeu_ps(dst + 28, v_a1); + v_float32x4 v_h; + v_float32x4 v_l; + v_float32x4 v_s; + v_load_deinterleave(src + i, v_h, v_l, v_s); + process(v_h, v_l, v_s); + v_float32x4 v_a = v_setall_f32(alpha); + v_store_interleave(dst, v_s, v_l, v_h, v_a); } - else + } else { + for (; i <= n - 12; i += 12, dst += dcn * 4) { - _mm_interleave_ps(v_h0, v_h1, v_l0, v_l1, v_s0, v_s1, v_a0, v_a1); - - _mm_storeu_ps(dst + 0, v_h0); - _mm_storeu_ps(dst + 4, v_h1); - _mm_storeu_ps(dst + 8, v_l0); - _mm_storeu_ps(dst + 12, v_l1); - _mm_storeu_ps(dst + 16, v_s0); - _mm_storeu_ps(dst + 20, v_s1); - _mm_storeu_ps(dst + 24, v_a0); - _mm_storeu_ps(dst + 28, v_a1); + v_float32x4 v_h; + v_float32x4 v_l; + v_float32x4 v_s; + v_load_deinterleave(src + i, v_h, v_l, v_s); + process(v_h, v_l, v_s); + v_float32x4 v_a = v_setall_f32(alpha); + v_store_interleave(dst, v_h, v_l, v_s, v_a); } } } } #endif + for( ; i < n; i += 3, dst += dcn ) { float h = src[i], l = src[i+1], s = src[i+2]; @@ -1050,7 +994,7 @@ struct HLS2RGB_f float p2 = l <= 0.5f ? l*(1 + s) : l + s - l*s; float p1 = 2*l - p2; - h *= _hscale; + h *= hscale; if( h < 0 ) do h += 6; while( h < 0 ); else if( h >= 6 ) @@ -1080,8 +1024,8 @@ struct HLS2RGB_f int dstcn, blueIdx; float hscale; - #if CV_SSE2 - bool haveSIMD; + #if CV_SIMD128 + bool hasSIMD; #endif }; From ecb4ea8f7677a5300b8dd767744ac4eddd5b82e9 Mon Sep 17 00:00:00 2001 From: k-shinotsuka Date: Sat, 28 Apr 2018 21:09:38 +0900 Subject: [PATCH 02/14] add universal intrinsics for RGB2HSV_f --- modules/imgproc/src/color_hsv.cpp | 91 ++++++++++++++++++++++++++++++- 1 file changed, 88 insertions(+), 3 deletions(-) diff --git a/modules/imgproc/src/color_hsv.cpp b/modules/imgproc/src/color_hsv.cpp index d8074b4fa5..2b01755c61 100644 --- a/modules/imgproc/src/color_hsv.cpp +++ b/modules/imgproc/src/color_hsv.cpp @@ -84,15 +84,97 @@ struct RGB2HSV_f typedef float channel_type; RGB2HSV_f(int _srccn, int _blueIdx, float _hrange) - : srccn(_srccn), blueIdx(_blueIdx), hrange(_hrange) {} + : srccn(_srccn), blueIdx(_blueIdx), hrange(_hrange) { + #if CV_SIMD128 + hasSIMD = hasSIMD128(); + #endif + } + + #if CV_SIMD128 + inline void process(v_float32x4& v_r, v_float32x4& v_g, + v_float32x4& v_b, float hscale) const + { + v_float32x4 v_min_rgb = v_min(v_min(v_r, v_g), v_b); + v_float32x4 v_max_rgb = v_max(v_max(v_r, v_g), v_b); + + v_float32x4 v_eps = v_setall_f32(FLT_EPSILON); + v_float32x4 v_diff = v_max_rgb - v_min_rgb; + v_float32x4 v_s = v_diff / (v_abs(v_max_rgb) + v_eps); + + v_float32x4 v_r_eq_max = v_r == v_max_rgb; + v_float32x4 v_g_eq_max = v_g == v_max_rgb; + v_float32x4 v_h = v_select(v_r_eq_max, v_g - v_b, + v_select(v_g_eq_max, v_b - v_r, v_r - v_g)); + v_float32x4 v_res = v_select(v_r_eq_max, (v_g < v_b) & v_setall_f32(360.0f), + v_select(v_g_eq_max, v_setall_f32(120.0f), v_setall_f32(240.0f))); + v_float32x4 v_rev_diff = v_setall_f32(60.0f) / (v_diff + v_eps); + v_r = v_muladd(v_h, v_rev_diff, v_res) * v_setall_f32(hscale); + + v_g = v_s; + v_b = v_max_rgb; + } + #endif void operator()(const float* src, float* dst, int n) const { - int i, bidx = blueIdx, scn = srccn; + int i = 0, bidx = blueIdx, scn = srccn; float hscale = hrange*(1.f/360.f); n *= 3; - for( i = 0; i < n; i += 3, src += scn ) + #if CV_SIMD128 + if (hasSIMD) + { + if (scn == 3) { + if (bidx) { + for ( ; i <= n - 12; i += 12, src += scn * 4) + { + v_float32x4 v_r; + v_float32x4 v_g; + v_float32x4 v_b; + v_load_deinterleave(src, v_r, v_g, v_b); + process(v_r, v_g, v_b, hscale); + v_store_interleave(dst + i, v_r, v_g, v_b); + } + } else { + for ( ; i <= n - 12; i += 12, src += scn * 4) + { + v_float32x4 v_r; + v_float32x4 v_g; + v_float32x4 v_b; + v_load_deinterleave(src, v_r, v_g, v_b); + process(v_b, v_g, v_r, hscale); + v_store_interleave(dst + i, v_b, v_g, v_r); + } + } + } else { // scn == 4 + if (bidx) { + for ( ; i <= n - 12; i += 12, src += scn * 4) + { + v_float32x4 v_r; + v_float32x4 v_g; + v_float32x4 v_b; + v_float32x4 v_a; + v_load_deinterleave(src, v_r, v_g, v_b, v_a); + process(v_r, v_g, v_b, hscale); + v_store_interleave(dst + i, v_r, v_g, v_b); + } + } else { + for ( ; i <= n - 12; i += 12, src += scn * 4) + { + v_float32x4 v_r; + v_float32x4 v_g; + v_float32x4 v_b; + v_float32x4 v_a; + v_load_deinterleave(src, v_r, v_g, v_b, v_a); + process(v_b, v_g, v_r, hscale); + v_store_interleave(dst + i, v_b, v_g, v_r); + } + } + } + } + #endif + + for( ; i < n; i += 3, src += scn ) { float b = src[bidx], g = src[1], r = src[bidx^2]; float h, s, v; @@ -125,6 +207,9 @@ struct RGB2HSV_f int srccn, blueIdx; float hrange; + #if CV_SIMD128 + bool hasSIMD; + #endif }; From e00a58a05f9895008e1a3d7939a125d5c3c3ba04 Mon Sep 17 00:00:00 2001 From: Sayed Adel Date: Thu, 3 May 2018 20:12:24 +0000 Subject: [PATCH 03/14] samples: add C++11 override to virtual methods --- .../cpp/tutorial_code/dnn/custom_layers.cpp | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/samples/cpp/tutorial_code/dnn/custom_layers.cpp b/samples/cpp/tutorial_code/dnn/custom_layers.cpp index 3fc9e61279..217e53659f 100644 --- a/samples/cpp/tutorial_code/dnn/custom_layers.cpp +++ b/samples/cpp/tutorial_code/dnn/custom_layers.cpp @@ -16,18 +16,18 @@ public: virtual bool getMemoryShapes(const std::vector > &inputs, const int requiredOutputs, std::vector > &outputs, - std::vector > &internals) const; + std::vector > &internals) const CV_OVERRIDE; //! [MyLayer::getMemoryShapes] //! [MyLayer::forward] - virtual void forward(std::vector &inputs, std::vector &outputs, std::vector &internals); + virtual void forward(std::vector &inputs, std::vector &outputs, std::vector &internals) CV_OVERRIDE; //! [MyLayer::forward] //! [MyLayer::finalize] - virtual void finalize(const std::vector &inputs, std::vector &outputs); + virtual void finalize(const std::vector &inputs, std::vector &outputs) CV_OVERRIDE; //! [MyLayer::finalize] - virtual void forward(cv::InputArrayOfArrays inputs, cv::OutputArrayOfArrays outputs, cv::OutputArrayOfArrays internals); + virtual void forward(cv::InputArrayOfArrays inputs, cv::OutputArrayOfArrays outputs, cv::OutputArrayOfArrays internals) CV_OVERRIDE; }; //! [A custom layer interface] @@ -49,7 +49,7 @@ public: virtual bool getMemoryShapes(const std::vector > &inputs, const int requiredOutputs, std::vector > &outputs, - std::vector > &internals) const + std::vector > &internals) const CV_OVERRIDE { CV_UNUSED(requiredOutputs); CV_UNUSED(internals); std::vector outShape(4); @@ -62,7 +62,7 @@ public: } // Implementation of this custom layer is based on https://github.com/cdmh/deeplab-public/blob/master/src/caffe/layers/interp_layer.cpp - virtual void forward(std::vector &inputs, std::vector &outputs, std::vector &internals) + virtual void forward(std::vector &inputs, std::vector &outputs, std::vector &internals) CV_OVERRIDE { CV_UNUSED(internals); cv::Mat& inp = *inputs[0]; @@ -105,7 +105,7 @@ public: } } - virtual void forward(cv::InputArrayOfArrays, cv::OutputArrayOfArrays, cv::OutputArrayOfArrays) {} + virtual void forward(cv::InputArrayOfArrays, cv::OutputArrayOfArrays, cv::OutputArrayOfArrays) CV_OVERRIDE {} private: int outWidth, outHeight; @@ -132,7 +132,7 @@ public: virtual bool getMemoryShapes(const std::vector > &inputs, const int requiredOutputs, std::vector > &outputs, - std::vector > &internals) const + std::vector > &internals) const CV_OVERRIDE { CV_UNUSED(requiredOutputs); CV_UNUSED(internals); std::vector outShape(4); @@ -146,7 +146,7 @@ public: // This implementation is based on a reference implementation from // https://github.com/tensorflow/tensorflow/blob/master/tensorflow/contrib/lite/kernels/internal/reference/reference_ops.h - virtual void forward(std::vector &inputs, std::vector &outputs, std::vector &internals) + virtual void forward(std::vector &inputs, std::vector &outputs, std::vector &internals) CV_OVERRIDE { CV_UNUSED(internals); cv::Mat& inp = *inputs[0]; @@ -187,7 +187,7 @@ public: } } - virtual void forward(cv::InputArrayOfArrays, cv::OutputArrayOfArrays, cv::OutputArrayOfArrays) {} + virtual void forward(cv::InputArrayOfArrays, cv::OutputArrayOfArrays, cv::OutputArrayOfArrays) CV_OVERRIDE {} private: static inline int offset(const cv::MatSize& size, int c, int x, int y, int b) From ed19da21abfade7e508ef18809d7424ff8189d13 Mon Sep 17 00:00:00 2001 From: Sayed Adel Date: Fri, 4 May 2018 23:09:38 +0000 Subject: [PATCH 04/14] core:ppc Several improvements on VSX(2) * add v_float64x2 support to v_rotate_* * treat float vector in v_check_any, vec_any_lt as int vector * add test case for v_rotate_left --- .../include/opencv2/core/hal/intrin_vsx.hpp | 37 ++++++++++++------- modules/core/test/test_intrin_utils.hpp | 13 ++++++- 2 files changed, 35 insertions(+), 15 deletions(-) diff --git a/modules/core/include/opencv2/core/hal/intrin_vsx.hpp b/modules/core/include/opencv2/core/hal/intrin_vsx.hpp index 3fff6651e3..8b76dd8487 100644 --- a/modules/core/include/opencv2/core/hal/intrin_vsx.hpp +++ b/modules/core/include/opencv2/core/hal/intrin_vsx.hpp @@ -589,7 +589,7 @@ inline _Tpvec v_rotate_left(const _Tpvec& a, const _Tpvec& b) return _Tpvec(vec_sld(a.val, b.val, CV_SHIFT)); } -#define OPENCV_IMPL_VSX_ROTATE_64(_Tpvec, suffix, rg1, rg2) \ +#define OPENCV_IMPL_VSX_ROTATE_64_2RG(_Tpvec, suffix, rg1, rg2) \ template \ inline _Tpvec v_rotate_##suffix(const _Tpvec& a, const _Tpvec& b) \ { \ @@ -598,11 +598,13 @@ inline _Tpvec v_rotate_##suffix(const _Tpvec& a, const _Tpvec& b) \ return imm ? b : a; \ } -OPENCV_IMPL_VSX_ROTATE_64(v_int64x2, right, a, b) -OPENCV_IMPL_VSX_ROTATE_64(v_uint64x2, right, a, b) +#define OPENCV_IMPL_VSX_ROTATE_64_2RG_LR(_Tpvec) \ +OPENCV_IMPL_VSX_ROTATE_64_2RG(_Tpvec, left, b, a) \ +OPENCV_IMPL_VSX_ROTATE_64_2RG(_Tpvec, right, a, b) -OPENCV_IMPL_VSX_ROTATE_64(v_int64x2, left, b, a) -OPENCV_IMPL_VSX_ROTATE_64(v_uint64x2, left, b, a) +OPENCV_IMPL_VSX_ROTATE_64_2RG_LR(v_float64x2) +OPENCV_IMPL_VSX_ROTATE_64_2RG_LR(v_uint64x2) +OPENCV_IMPL_VSX_ROTATE_64_2RG_LR(v_int64x2) /* Extract */ template @@ -716,26 +718,33 @@ inline int v_signmask(const v_uint64x2& a) inline int v_signmask(const v_float64x2& a) { return v_signmask(v_reinterpret_as_s64(a)); } - template inline bool v_check_all(const _Tpvec& a) -{ return vec_all_lt(a.val, _Tpvec().val);} -inline bool v_check_all(const v_uint8x16 &a) +{ return vec_all_lt(a.val, _Tpvec().val); } +inline bool v_check_all(const v_uint8x16& a) { return v_check_all(v_reinterpret_as_s8(a)); } -inline bool v_check_all(const v_uint16x8 &a) +inline bool v_check_all(const v_uint16x8& a) { return v_check_all(v_reinterpret_as_s16(a)); } -inline bool v_check_all(const v_uint32x4 &a) +inline bool v_check_all(const v_uint32x4& a) { return v_check_all(v_reinterpret_as_s32(a)); } +inline bool v_check_all(const v_float32x4& a) +{ return v_check_all(v_reinterpret_as_s32(a)); } +inline bool v_check_all(const v_float64x2& a) +{ return v_check_all(v_reinterpret_as_s64(a)); } template inline bool v_check_any(const _Tpvec& a) -{ return vec_any_lt(a.val, _Tpvec().val);} -inline bool v_check_any(const v_uint8x16 &a) +{ return vec_any_lt(a.val, _Tpvec().val); } +inline bool v_check_any(const v_uint8x16& a) { return v_check_any(v_reinterpret_as_s8(a)); } -inline bool v_check_any(const v_uint16x8 &a) +inline bool v_check_any(const v_uint16x8& a) { return v_check_any(v_reinterpret_as_s16(a)); } -inline bool v_check_any(const v_uint32x4 &a) +inline bool v_check_any(const v_uint32x4& a) { return v_check_any(v_reinterpret_as_s32(a)); } +inline bool v_check_any(const v_float32x4& a) +{ return v_check_any(v_reinterpret_as_s32(a)); } +inline bool v_check_any(const v_float64x2& a) +{ return v_check_any(v_reinterpret_as_s64(a)); } ////////// Other math ///////// diff --git a/modules/core/test/test_intrin_utils.hpp b/modules/core/test/test_intrin_utils.hpp index 7a21c9eb56..7579d9cf05 100644 --- a/modules/core/test/test_intrin_utils.hpp +++ b/modules/core/test/test_intrin_utils.hpp @@ -837,17 +837,28 @@ template struct TheTest Data resC = v_rotate_right(a); Data resD = v_rotate_right(a, b); + Data resE = v_rotate_left(a); + Data resF = v_rotate_left(a, b); + for (int i = 0; i < R::nlanes; ++i) { if (i + s >= R::nlanes) { EXPECT_EQ((LaneType)0, resC[i]); EXPECT_EQ(dataB[i - R::nlanes + s], resD[i]); + + EXPECT_EQ((LaneType)0, resE[i - R::nlanes + s]); + EXPECT_EQ(dataB[i], resF[i - R::nlanes + s]); } else + { EXPECT_EQ(dataA[i + s], resC[i]); - } + EXPECT_EQ(dataA[i + s], resD[i]); + EXPECT_EQ(dataA[i], resE[i + s]); + EXPECT_EQ(dataA[i], resF[i + s]); + } + } return *this; } From a31aff452a4b088e026678202a659246ceefc9ea Mon Sep 17 00:00:00 2001 From: abhi-jha Date: Sat, 5 May 2018 00:56:29 +0530 Subject: [PATCH 05/14] fix spelling error --- .../interoperability_with_OpenCV_1.markdown | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/tutorials/core/interoperability_with_OpenCV_1/interoperability_with_OpenCV_1.markdown b/doc/tutorials/core/interoperability_with_OpenCV_1/interoperability_with_OpenCV_1.markdown index 42fdb3483e..298c45b3b1 100644 --- a/doc/tutorials/core/interoperability_with_OpenCV_1/interoperability_with_OpenCV_1.markdown +++ b/doc/tutorials/core/interoperability_with_OpenCV_1/interoperability_with_OpenCV_1.markdown @@ -69,7 +69,7 @@ CvMat* mI = &I.operator CvMat(); @endcode One of the biggest complaints of the C interface is that it leaves all the memory management to you. You need to figure out when it is safe to release your unused objects and make sure you do so before -the program finishes or you could have troublesome memory leeks. To work around this issue in OpenCV +the program finishes or you could have troublesome memory leaks. To work around this issue in OpenCV there is introduced a sort of smart pointer. This will automatically release the object when it's no longer in use. To use this declare the pointers as a specialization of the *Ptr* : @code{.cpp} From f2fbfd7a10468946199c43c4e446d03505c23ac0 Mon Sep 17 00:00:00 2001 From: cDc Date: Tue, 17 Apr 2018 13:30:56 +0300 Subject: [PATCH 06/14] fix bug in Matx::inv() --- modules/core/include/opencv2/core/operations.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/core/include/opencv2/core/operations.hpp b/modules/core/include/opencv2/core/operations.hpp index e5d1a7e80a..75864ea822 100644 --- a/modules/core/include/opencv2/core/operations.hpp +++ b/modules/core/include/opencv2/core/operations.hpp @@ -194,8 +194,8 @@ Matx<_Tp, n, m> Matx<_Tp, m, n>::inv(int method, bool *p_is_ok /*= NULL*/) const { Matx<_Tp, n, m> b; bool ok; - if( method == DECOMP_LU || method == DECOMP_CHOLESKY ) - ok = cv::internal::Matx_FastInvOp<_Tp, m>()(*this, b, method); + if( m == n && (method == DECOMP_LU || method == DECOMP_CHOLESKY) ) + ok = cv::internal::Matx_FastInvOp<_Tp, m>()(*reinterpret_cast*>(this), reinterpret_cast&>(b), method); else { Mat A(*this, false), B(b, false); From 115b10ed32c1c9197ba58da5beb9a3386e2323ea Mon Sep 17 00:00:00 2001 From: Sayed Adel Date: Sat, 5 May 2018 09:38:55 +0000 Subject: [PATCH 07/14] photo:ppc fix MergeRobertson test --- modules/photo/test/test_hdr.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/photo/test/test_hdr.cpp b/modules/photo/test/test_hdr.cpp index 9e2aeaa94f..c4bf536278 100644 --- a/modules/photo/test/test_hdr.cpp +++ b/modules/photo/test/test_hdr.cpp @@ -213,7 +213,7 @@ TEST(Photo_MergeRobertson, regression) loadImage(test_path + "merge/robertson.hdr", expected); merge->process(images, result, times); -#ifdef __aarch64__ +#if defined(__aarch64__) || defined(__PPC64__) const float eps = 6.f; #else const float eps = 5.f; From a58c9d4d630926ed81beda1a8f7f8ec52165949e Mon Sep 17 00:00:00 2001 From: Tomoaki Teshima Date: Tue, 8 May 2018 00:35:18 +0900 Subject: [PATCH 08/14] arm: fix build error of v_rotate_left * remove meaningless tests --- modules/core/test/test_intrin.cpp | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/modules/core/test/test_intrin.cpp b/modules/core/test/test_intrin.cpp index 4171babc03..4cf098bc0c 100644 --- a/modules/core/test/test_intrin.cpp +++ b/modules/core/test/test_intrin.cpp @@ -35,7 +35,7 @@ TEST(hal_intrin, uint8x16) { .test_pack_u<1>().test_pack_u<2>().test_pack_u<3>().test_pack_u<8>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<8>().test_extract<15>() - .test_rotate<0>().test_rotate<1>().test_rotate<8>().test_rotate<15>() + .test_rotate<1>().test_rotate<8>().test_rotate<15>() ; } @@ -57,7 +57,7 @@ TEST(hal_intrin, int8x16) { .test_pack<1>().test_pack<2>().test_pack<3>().test_pack<8>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<8>().test_extract<15>() - .test_rotate<0>().test_rotate<1>().test_rotate<8>().test_rotate<15>() + .test_rotate<1>().test_rotate<8>().test_rotate<15>() ; } @@ -85,7 +85,7 @@ TEST(hal_intrin, uint16x8) { .test_pack_u<1>().test_pack_u<2>().test_pack_u<7>().test_pack_u<16>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<4>().test_extract<7>() - .test_rotate<0>().test_rotate<1>().test_rotate<4>().test_rotate<7>() + .test_rotate<1>().test_rotate<4>().test_rotate<7>() ; } @@ -112,7 +112,7 @@ TEST(hal_intrin, int16x8) { .test_pack<1>().test_pack<2>().test_pack<7>().test_pack<16>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<4>().test_extract<7>() - .test_rotate<0>().test_rotate<1>().test_rotate<4>().test_rotate<7>() + .test_rotate<1>().test_rotate<4>().test_rotate<7>() ; } @@ -138,7 +138,7 @@ TEST(hal_intrin, uint32x4) { .test_pack<1>().test_pack<2>().test_pack<15>().test_pack<32>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<2>().test_extract<3>() - .test_rotate<0>().test_rotate<1>().test_rotate<2>().test_rotate<3>() + .test_rotate<1>().test_rotate<2>().test_rotate<3>() .test_transpose() ; } @@ -162,7 +162,7 @@ TEST(hal_intrin, int32x4) { .test_pack<1>().test_pack<2>().test_pack<15>().test_pack<32>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<2>().test_extract<3>() - .test_rotate<0>().test_rotate<1>().test_rotate<2>().test_rotate<3>() + .test_rotate<1>().test_rotate<2>().test_rotate<3>() .test_float_cvt32() .test_float_cvt64() .test_transpose() @@ -178,7 +178,7 @@ TEST(hal_intrin, uint64x2) { .test_shift<1>().test_shift<8>() .test_logic() .test_extract<0>().test_extract<1>() - .test_rotate<0>().test_rotate<1>() + .test_rotate<1>() ; } @@ -189,7 +189,7 @@ TEST(hal_intrin, int64x2) { .test_shift<1>().test_shift<8>() .test_logic() .test_extract<0>().test_extract<1>() - .test_rotate<0>().test_rotate<1>() + .test_rotate<1>() ; } @@ -216,7 +216,7 @@ TEST(hal_intrin, float32x4) { .test_transpose() .test_reduce_sum4() .test_extract<0>().test_extract<1>().test_extract<2>().test_extract<3>() - .test_rotate<0>().test_rotate<1>().test_rotate<2>().test_rotate<3>() + .test_rotate<1>().test_rotate<2>().test_rotate<3>() ; } @@ -236,7 +236,7 @@ TEST(hal_intrin, float64x2) { .test_float_math() .test_float_cvt32() .test_extract<0>().test_extract<1>() - .test_rotate<0>().test_rotate<1>() + .test_rotate<1>() ; } #endif From ba000bf8ea9433b7899f712d9f6fbe52c6937322 Mon Sep 17 00:00:00 2001 From: Damien Picard Date: Tue, 8 May 2018 15:02:01 +0200 Subject: [PATCH 09/14] Fixes 11480 : allow to force the video codec used to read videos with the ffmpeg backend --- modules/videoio/src/cap_ffmpeg_impl.hpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/modules/videoio/src/cap_ffmpeg_impl.hpp b/modules/videoio/src/cap_ffmpeg_impl.hpp index de600302a1..0317831d3a 100644 --- a/modules/videoio/src/cap_ffmpeg_impl.hpp +++ b/modules/videoio/src/cap_ffmpeg_impl.hpp @@ -870,7 +870,12 @@ bool CvCapture_FFMPEG::open( const char* _filename ) int enc_width = enc->width; int enc_height = enc->height; - AVCodec *codec = avcodec_find_decoder(enc->codec_id); + AVCodec *codec; + if(av_dict_get(dict, "video_codec", NULL, 0) == NULL) { + codec = avcodec_find_decoder(enc->codec_id); + } else { + codec = avcodec_find_decoder_by_name(av_dict_get(dict, "video_codec", NULL, 0)->value); + } if (!codec || #if LIBAVCODEC_VERSION_INT >= ((53<<16)+(8<<8)+0) avcodec_open2(enc, codec, NULL) From f7ea6b12ea42a99c90537b434c85d5f08c1cd54b Mon Sep 17 00:00:00 2001 From: Tomoaki Teshima Date: Wed, 9 May 2018 00:23:49 +0900 Subject: [PATCH 10/14] cuda: fix build error of sample --- modules/stitching/include/opencv2/stitching/detail/blenders.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/modules/stitching/include/opencv2/stitching/detail/blenders.hpp b/modules/stitching/include/opencv2/stitching/detail/blenders.hpp index 4c14340722..542f1e42fb 100644 --- a/modules/stitching/include/opencv2/stitching/detail/blenders.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/blenders.hpp @@ -48,6 +48,7 @@ #endif #include "opencv2/core.hpp" +#include "opencv2/core/cuda.hpp" namespace cv { namespace detail { From 70455a959dcc2b4511fb903775521d27b8a20938 Mon Sep 17 00:00:00 2001 From: shengyu Date: Wed, 9 May 2018 21:15:51 +0800 Subject: [PATCH 11/14] tab to space --- modules/videoio/src/cap_qt.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/modules/videoio/src/cap_qt.cpp b/modules/videoio/src/cap_qt.cpp index a75ff2ac16..cb416e6849 100644 --- a/modules/videoio/src/cap_qt.cpp +++ b/modules/videoio/src/cap_qt.cpp @@ -634,11 +634,11 @@ static int icvOpenCamera_QT (CvCapture_QT_Cam * capture, const int index) OPENCV_ASSERT (capture, "icvOpenCamera_QT", "'capture' is a NULL-pointer"); OPENCV_ASSERT (index >=0, "icvOpenCamera_QT", "camera index is negative"); - ComponentDescription component_description; - Component component = 0; + ComponentDescription component_description; + Component component = 0; int number_of_inputs = 0; Rect myRect; - ComponentResult result = noErr; + ComponentResult result = noErr; // travers all components and count video digitizer channels From 1b3e0783f47c430f6cd02cc099ecb49ffa01641b Mon Sep 17 00:00:00 2001 From: cabelo Date: Tue, 8 May 2018 01:07:23 -0300 Subject: [PATCH 12/14] select the device (video capture) --- doc/tutorials/dnn/dnn_yolo/dnn_yolo.markdown | 4 ++-- samples/dnn/object_detection.cpp | 9 +++++---- samples/dnn/openpose.cpp | 8 ++++---- samples/dnn/segmentation.cpp | 9 +++++---- 4 files changed, 16 insertions(+), 14 deletions(-) diff --git a/doc/tutorials/dnn/dnn_yolo/dnn_yolo.markdown b/doc/tutorials/dnn/dnn_yolo/dnn_yolo.markdown index e9f446db36..968b3faaca 100644 --- a/doc/tutorials/dnn/dnn_yolo/dnn_yolo.markdown +++ b/doc/tutorials/dnn/dnn_yolo/dnn_yolo.markdown @@ -29,7 +29,7 @@ Execute in webcam: @code{.bash} -$ example_dnn_object_detection --config=[PATH-TO-DARKNET]/cfg/yolo.cfg --model=[PATH-TO-DARKNET]/yolo.weights --classes=object_detection_classes_pascal_voc.txt --width=416 --height=416 --scale=0.00392 +$ example_dnn_object_detection --config=[PATH-TO-DARKNET]/cfg/yolo.cfg --model=[PATH-TO-DARKNET]/yolo.weights --classes=object_detection_classes_pascal_voc.txt --width=416 --height=416 --scale=0.00392 --rgb @endcode @@ -37,7 +37,7 @@ Execute with image or video file: @code{.bash} -$ example_dnn_object_detection --config=[PATH-TO-DARKNET]/cfg/yolo.cfg --model=[PATH-TO-DARKNET]/yolo.weights --classes=object_detection_classes_pascal_voc.txt --width=416 --height=416 --scale=0.00392 --input=[PATH-TO-IMAGE-OR-VIDEO-FILE] +$ example_dnn_object_detection --config=[PATH-TO-DARKNET]/cfg/yolo.cfg --model=[PATH-TO-DARKNET]/yolo.weights --classes=object_detection_classes_pascal_voc.txt --width=416 --height=416 --scale=0.00392 --input=[PATH-TO-IMAGE-OR-VIDEO-FILE] --rgb @endcode diff --git a/samples/dnn/object_detection.cpp b/samples/dnn/object_detection.cpp index 5ff537bdbd..1298d7e39e 100644 --- a/samples/dnn/object_detection.cpp +++ b/samples/dnn/object_detection.cpp @@ -7,12 +7,13 @@ const char* keys = "{ help h | | Print help message. }" - "{ input i | | Path to input image or video file. Skip this argument to capture frames from a camera.}" + "{ device | 0 | camera device number. }" + "{ input i | | Path to input image or video file. Skip this argument to capture frames from a camera. }" "{ model m | | Path to a binary file of model contains trained weights. " "It could be a file with extensions .caffemodel (Caffe), " - ".pb (TensorFlow), .t7 or .net (Torch), .weights (Darknet) }" + ".pb (TensorFlow), .t7 or .net (Torch), .weights (Darknet).}" "{ config c | | Path to a text file of model contains network configuration. " - "It could be a file with extensions .prototxt (Caffe), .pbtxt (TensorFlow), .cfg (Darknet) }" + "It could be a file with extensions .prototxt (Caffe), .pbtxt (TensorFlow), .cfg (Darknet).}" "{ framework f | | Optional name of an origin framework of the model. Detect it automatically if it does not set. }" "{ classes | | Optional path to a text file with names of classes to label detected objects. }" "{ mean | | Preprocess input image by subtracting mean values. Mean values should be in BGR order and delimited by spaces. }" @@ -91,7 +92,7 @@ int main(int argc, char** argv) if (parser.has("input")) cap.open(parser.get("input")); else - cap.open(0); + cap.open(parser.get("device")); // Process frames. Mat frame, blob; diff --git a/samples/dnn/openpose.cpp b/samples/dnn/openpose.cpp index bc95c60023..da9315426a 100644 --- a/samples/dnn/openpose.cpp +++ b/samples/dnn/openpose.cpp @@ -61,12 +61,16 @@ int main(int argc, char **argv) "{ p proto | | (required) model configuration, e.g. hand/pose.prototxt }" "{ m model | | (required) model weights, e.g. hand/pose_iter_102000.caffemodel }" "{ i image | | (required) path to image file (containing a single person, or hand) }" + "{ width | 368 | Preprocess input image by resizing to a specific width. }" + "{ height | 368 | Preprocess input image by resizing to a specific height. }" "{ t threshold | 0.1 | threshold or confidence value for the heatmap }" ); String modelTxt = parser.get("proto"); String modelBin = parser.get("model"); String imageFile = parser.get("image"); + int W_in = parser.get("width"); + int H_in = parser.get("height"); float thresh = parser.get("threshold"); if (parser.get("help") || modelTxt.empty() || modelBin.empty() || imageFile.empty()) { @@ -75,10 +79,6 @@ int main(int argc, char **argv) return 0; } - // fixed input size for the pretrained network - int W_in = 368; - int H_in = 368; - // read the network model Net net = readNetFromCaffe(modelTxt, modelBin); diff --git a/samples/dnn/segmentation.cpp b/samples/dnn/segmentation.cpp index a0eb15bc86..252140a275 100644 --- a/samples/dnn/segmentation.cpp +++ b/samples/dnn/segmentation.cpp @@ -7,12 +7,13 @@ const char* keys = "{ help h | | Print help message. }" - "{ input i | | Path to input image or video file. Skip this argument to capture frames from a camera.}" + "{ device | 0 | camera device number. }" + "{ input i | | Path to input image or video file. Skip this argument to capture frames from a camera. }" "{ model m | | Path to a binary file of model contains trained weights. " "It could be a file with extensions .caffemodel (Caffe), " - ".pb (TensorFlow), .t7 or .net (Torch), .weights (Darknet) }" + ".pb (TensorFlow), .t7 or .net (Torch), .weights (Darknet). }" "{ config c | | Path to a text file of model contains network configuration. " - "It could be a file with extensions .prototxt (Caffe), .pbtxt (TensorFlow), .cfg (Darknet) }" + "It could be a file with extensions .prototxt (Caffe), .pbtxt (TensorFlow), .cfg (Darknet). }" "{ framework f | | Optional name of an origin framework of the model. Detect it automatically if it does not set. }" "{ classes | | Optional path to a text file with names of classes. }" "{ colors | | Optional path to a text file with colors for an every class. " @@ -111,7 +112,7 @@ int main(int argc, char** argv) if (parser.has("input")) cap.open(parser.get("input")); else - cap.open(0); + cap.open(parser.get("device")); //! [Open a video file or an image file or a camera stream] // Process frames. From 352510cc19bab69badcfdcd49c1893437dc9ecdd Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Mon, 7 May 2018 02:54:10 +0300 Subject: [PATCH 13/14] core: fix ARM intrinsincs '0' is specific case (make no sence as a standalone operation), but it can be useful in template-based programming. reverts commit: a58c9d4d630926ed81beda1a8f7f8ec52165949e --- .../include/opencv2/core/hal/intrin_neon.hpp | 6 +++++- modules/core/test/test_intrin.cpp | 20 +++++++++---------- 2 files changed, 15 insertions(+), 11 deletions(-) diff --git a/modules/core/include/opencv2/core/hal/intrin_neon.hpp b/modules/core/include/opencv2/core/hal/intrin_neon.hpp index 95c9bfb1fe..9dadab57ea 100644 --- a/modules/core/include/opencv2/core/hal/intrin_neon.hpp +++ b/modules/core/include/opencv2/core/hal/intrin_neon.hpp @@ -786,10 +786,14 @@ template inline _Tpvec v_rotate_right(const _Tpvec& a) \ { return _Tpvec(vextq_##suffix(a.val, vdupq_n_##suffix(0), n)); } \ template inline _Tpvec v_rotate_left(const _Tpvec& a) \ { return _Tpvec(vextq_##suffix(vdupq_n_##suffix(0), a.val, _Tpvec::nlanes - n)); } \ +template<> inline _Tpvec v_rotate_left<0>(const _Tpvec& a) \ +{ return a; } \ template inline _Tpvec v_rotate_right(const _Tpvec& a, const _Tpvec& b) \ { return _Tpvec(vextq_##suffix(a.val, b.val, n)); } \ template inline _Tpvec v_rotate_left(const _Tpvec& a, const _Tpvec& b) \ -{ return _Tpvec(vextq_##suffix(b.val, a.val, _Tpvec::nlanes - n)); } +{ return _Tpvec(vextq_##suffix(b.val, a.val, _Tpvec::nlanes - n)); } \ +template<> inline _Tpvec v_rotate_left<0>(const _Tpvec& a, const _Tpvec& b) \ +{ CV_UNUSED(b); return a; } OPENCV_HAL_IMPL_NEON_ROTATE_OP(v_uint8x16, u8) OPENCV_HAL_IMPL_NEON_ROTATE_OP(v_int8x16, s8) diff --git a/modules/core/test/test_intrin.cpp b/modules/core/test/test_intrin.cpp index 4cf098bc0c..4171babc03 100644 --- a/modules/core/test/test_intrin.cpp +++ b/modules/core/test/test_intrin.cpp @@ -35,7 +35,7 @@ TEST(hal_intrin, uint8x16) { .test_pack_u<1>().test_pack_u<2>().test_pack_u<3>().test_pack_u<8>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<8>().test_extract<15>() - .test_rotate<1>().test_rotate<8>().test_rotate<15>() + .test_rotate<0>().test_rotate<1>().test_rotate<8>().test_rotate<15>() ; } @@ -57,7 +57,7 @@ TEST(hal_intrin, int8x16) { .test_pack<1>().test_pack<2>().test_pack<3>().test_pack<8>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<8>().test_extract<15>() - .test_rotate<1>().test_rotate<8>().test_rotate<15>() + .test_rotate<0>().test_rotate<1>().test_rotate<8>().test_rotate<15>() ; } @@ -85,7 +85,7 @@ TEST(hal_intrin, uint16x8) { .test_pack_u<1>().test_pack_u<2>().test_pack_u<7>().test_pack_u<16>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<4>().test_extract<7>() - .test_rotate<1>().test_rotate<4>().test_rotate<7>() + .test_rotate<0>().test_rotate<1>().test_rotate<4>().test_rotate<7>() ; } @@ -112,7 +112,7 @@ TEST(hal_intrin, int16x8) { .test_pack<1>().test_pack<2>().test_pack<7>().test_pack<16>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<4>().test_extract<7>() - .test_rotate<1>().test_rotate<4>().test_rotate<7>() + .test_rotate<0>().test_rotate<1>().test_rotate<4>().test_rotate<7>() ; } @@ -138,7 +138,7 @@ TEST(hal_intrin, uint32x4) { .test_pack<1>().test_pack<2>().test_pack<15>().test_pack<32>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<2>().test_extract<3>() - .test_rotate<1>().test_rotate<2>().test_rotate<3>() + .test_rotate<0>().test_rotate<1>().test_rotate<2>().test_rotate<3>() .test_transpose() ; } @@ -162,7 +162,7 @@ TEST(hal_intrin, int32x4) { .test_pack<1>().test_pack<2>().test_pack<15>().test_pack<32>() .test_unpack() .test_extract<0>().test_extract<1>().test_extract<2>().test_extract<3>() - .test_rotate<1>().test_rotate<2>().test_rotate<3>() + .test_rotate<0>().test_rotate<1>().test_rotate<2>().test_rotate<3>() .test_float_cvt32() .test_float_cvt64() .test_transpose() @@ -178,7 +178,7 @@ TEST(hal_intrin, uint64x2) { .test_shift<1>().test_shift<8>() .test_logic() .test_extract<0>().test_extract<1>() - .test_rotate<1>() + .test_rotate<0>().test_rotate<1>() ; } @@ -189,7 +189,7 @@ TEST(hal_intrin, int64x2) { .test_shift<1>().test_shift<8>() .test_logic() .test_extract<0>().test_extract<1>() - .test_rotate<1>() + .test_rotate<0>().test_rotate<1>() ; } @@ -216,7 +216,7 @@ TEST(hal_intrin, float32x4) { .test_transpose() .test_reduce_sum4() .test_extract<0>().test_extract<1>().test_extract<2>().test_extract<3>() - .test_rotate<1>().test_rotate<2>().test_rotate<3>() + .test_rotate<0>().test_rotate<1>().test_rotate<2>().test_rotate<3>() ; } @@ -236,7 +236,7 @@ TEST(hal_intrin, float64x2) { .test_float_math() .test_float_cvt32() .test_extract<0>().test_extract<1>() - .test_rotate<1>() + .test_rotate<0>().test_rotate<1>() ; } #endif From ed86bd34b190b2d6dba3f8bccd6070f0e5062b93 Mon Sep 17 00:00:00 2001 From: Namgoo Lee Date: Thu, 10 May 2018 05:44:34 +0900 Subject: [PATCH 14/14] Merge pull request #11483 from nglee:dev_cudaCannyStreamIssue cuda_canny : multi stream safety (#11483) * CUDA_ImgProc/Canny Asynchronous test * cuda_canny : multi stream safety (1/3) - Convert global variable canny::counter to class local variable * cuda_canny : multi stream safety (2/3) - Use texture objects rather than texture reference for cc >= 3.0, since texture reference must be declared as a static global variable which results in race condition when ran concurrently * cuda_canny : multi stream safety (3/3) - Refrain from using global variable in row_filter and column_filter (converts column_filter::c_kernel and row_filter::c_kernel to local variables) * Fixes #11193 --- .../cudafilters/src/cuda/column_filter.hpp | 19 +- modules/cudafilters/src/cuda/row_filter.hpp | 19 +- modules/cudaimgproc/src/canny.cpp | 17 +- modules/cudaimgproc/src/cuda/canny.cu | 232 +++++++++++++++--- modules/cudaimgproc/test/test_canny.cpp | 59 ++++- 5 files changed, 276 insertions(+), 70 deletions(-) diff --git a/modules/cudafilters/src/cuda/column_filter.hpp b/modules/cudafilters/src/cuda/column_filter.hpp index 7dc339ca91..e93fc836fa 100644 --- a/modules/cudafilters/src/cuda/column_filter.hpp +++ b/modules/cudafilters/src/cuda/column_filter.hpp @@ -52,10 +52,8 @@ namespace column_filter { #define MAX_KERNEL_SIZE 32 - __constant__ float c_kernel[MAX_KERNEL_SIZE]; - template - __global__ void linearColumnFilter(const PtrStepSz src, PtrStep dst, const int anchor, const B brd) + __global__ void linearColumnFilter(const PtrStepSz src, PtrStep dst, const float* kernel, const int anchor, const B brd) { #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200) const int BLOCK_DIM_X = 16; @@ -135,7 +133,7 @@ namespace column_filter #pragma unroll for (int k = 0; k < KSIZE; ++k) - sum = sum + smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y - anchor + k][threadIdx.x] * c_kernel[k]; + sum = sum + smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y - anchor + k][threadIdx.x] * kernel[k]; dst(y, x) = saturate_cast(sum); } @@ -143,7 +141,7 @@ namespace column_filter } template class B> - void caller(PtrStepSz src, PtrStepSz dst, int anchor, int cc, cudaStream_t stream) + void caller(PtrStepSz src, PtrStepSz dst, const float* kernel, int anchor, int cc, cudaStream_t stream) { int BLOCK_DIM_X; int BLOCK_DIM_Y; @@ -167,7 +165,7 @@ namespace column_filter B brd(src.rows); - linearColumnFilter<<>>(src, dst, anchor, brd); + linearColumnFilter<<>>(src, dst, kernel, anchor, brd); cudaSafeCall( cudaGetLastError() ); @@ -181,7 +179,7 @@ namespace filter template void linearColumn(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream) { - typedef void (*caller_t)(PtrStepSz src, PtrStepSz dst, int anchor, int cc, cudaStream_t stream); + typedef void (*caller_t)(PtrStepSz src, PtrStepSz dst, const float* kernel, int anchor, int cc, cudaStream_t stream); static const caller_t callers[5][33] = { @@ -362,11 +360,6 @@ namespace filter } }; - if (stream == 0) - cudaSafeCall( cudaMemcpyToSymbol(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); - else - cudaSafeCall( cudaMemcpyToSymbolAsync(column_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); - - callers[brd_type][ksize]((PtrStepSz)src, (PtrStepSz)dst, anchor, cc, stream); + callers[brd_type][ksize]((PtrStepSz)src, (PtrStepSz)dst, kernel, anchor, cc, stream); } } diff --git a/modules/cudafilters/src/cuda/row_filter.hpp b/modules/cudafilters/src/cuda/row_filter.hpp index 80eab59b44..4a4be36f9a 100644 --- a/modules/cudafilters/src/cuda/row_filter.hpp +++ b/modules/cudafilters/src/cuda/row_filter.hpp @@ -52,10 +52,8 @@ namespace row_filter { #define MAX_KERNEL_SIZE 32 - __constant__ float c_kernel[MAX_KERNEL_SIZE]; - template - __global__ void linearRowFilter(const PtrStepSz src, PtrStep dst, const int anchor, const B brd) + __global__ void linearRowFilter(const PtrStepSz src, PtrStep dst, const float* kernel, const int anchor, const B brd) { #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200) const int BLOCK_DIM_X = 32; @@ -135,7 +133,7 @@ namespace row_filter #pragma unroll for (int k = 0; k < KSIZE; ++k) - sum = sum + smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X - anchor + k] * c_kernel[k]; + sum = sum + smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X - anchor + k] * kernel[k]; dst(y, x) = saturate_cast(sum); } @@ -143,7 +141,7 @@ namespace row_filter } template class B> - void caller(PtrStepSz src, PtrStepSz dst, int anchor, int cc, cudaStream_t stream) + void caller(PtrStepSz src, PtrStepSz dst, const float* kernel, int anchor, int cc, cudaStream_t stream) { int BLOCK_DIM_X; int BLOCK_DIM_Y; @@ -167,7 +165,7 @@ namespace row_filter B brd(src.cols); - linearRowFilter<<>>(src, dst, anchor, brd); + linearRowFilter<<>>(src, dst, kernel, anchor, brd); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -180,7 +178,7 @@ namespace filter template void linearRow(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream) { - typedef void (*caller_t)(PtrStepSz src, PtrStepSz dst, int anchor, int cc, cudaStream_t stream); + typedef void (*caller_t)(PtrStepSz src, PtrStepSz dst, const float* kernel, int anchor, int cc, cudaStream_t stream); static const caller_t callers[5][33] = { @@ -361,11 +359,6 @@ namespace filter } }; - if (stream == 0) - cudaSafeCall( cudaMemcpyToSymbol(row_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); - else - cudaSafeCall( cudaMemcpyToSymbolAsync(row_filter::c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); - - callers[brd_type][ksize]((PtrStepSz)src, (PtrStepSz)dst, anchor, cc, stream); + callers[brd_type][ksize]((PtrStepSz)src, (PtrStepSz)dst, kernel, anchor, cc, stream); } } diff --git a/modules/cudaimgproc/src/canny.cpp b/modules/cudaimgproc/src/canny.cpp index 75e53cf3ec..8c3fd4a2b5 100644 --- a/modules/cudaimgproc/src/canny.cpp +++ b/modules/cudaimgproc/src/canny.cpp @@ -58,9 +58,9 @@ namespace canny void calcMap(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, PtrStepSzi map, float low_thresh, float high_thresh, cudaStream_t stream); - void edgesHysteresisLocal(PtrStepSzi map, short2* st1, cudaStream_t stream); + void edgesHysteresisLocal(PtrStepSzi map, short2* st1, int* d_counter, cudaStream_t stream); - void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2, cudaStream_t stream); + void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2, int* d_counter, cudaStream_t stream); void getEdges(PtrStepSzi map, PtrStepSzb dst, cudaStream_t stream); } @@ -127,6 +127,8 @@ namespace Ptr filterDX_, filterDY_; #endif int old_apperture_size_; + + int *d_counter; }; void CannyImpl::detect(InputArray _image, OutputArray _edges, Stream& stream) @@ -218,12 +220,17 @@ namespace void CannyImpl::CannyCaller(GpuMat& edges, Stream& stream) { - map_.setTo(Scalar::all(0)); + map_.setTo(Scalar::all(0), stream); + canny::calcMap(dx_, dy_, mag_, map_, static_cast(low_thresh_), static_cast(high_thresh_), StreamAccessor::getStream(stream)); - canny::edgesHysteresisLocal(map_, st1_.ptr(), StreamAccessor::getStream(stream)); + cudaSafeCall( cudaMalloc(&d_counter, sizeof(int)) ); - canny::edgesHysteresisGlobal(map_, st1_.ptr(), st2_.ptr(), StreamAccessor::getStream(stream)); + canny::edgesHysteresisLocal(map_, st1_.ptr(), d_counter, StreamAccessor::getStream(stream)); + + canny::edgesHysteresisGlobal(map_, st1_.ptr(), st2_.ptr(), d_counter, StreamAccessor::getStream(stream)); + + cudaSafeCall( cudaFree(d_counter) ); canny::getEdges(map_, edges, StreamAccessor::getStream(stream)); } diff --git a/modules/cudaimgproc/src/cuda/canny.cu b/modules/cudaimgproc/src/cuda/canny.cu index e0ba515693..4418b8e5eb 100644 --- a/modules/cudaimgproc/src/cuda/canny.cu +++ b/modules/cudaimgproc/src/cuda/canny.cu @@ -47,6 +47,7 @@ #include "opencv2/core/cuda/transform.hpp" #include "opencv2/core/cuda/functional.hpp" #include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda.hpp" using namespace cv::cuda; using namespace cv::cuda::device; @@ -102,6 +103,20 @@ namespace canny } }; + struct SrcTexObject + { + int xoff; + int yoff; + cudaTextureObject_t tex_src_object; + __host__ SrcTexObject(int _xoff, int _yoff, cudaTextureObject_t _tex_src_object) : xoff(_xoff), yoff(_yoff), tex_src_object(_tex_src_object) { } + + __device__ __forceinline__ int operator ()(int y, int x) const + { + return tex2D(tex_src_object, x + xoff, y + yoff); + } + + }; + template __global__ void calcMagnitudeKernel(const SrcTex src, PtrStepi dx, PtrStepi dy, PtrStepSzf mag, const Norm norm) { @@ -120,29 +135,94 @@ namespace canny mag(y, x) = norm(dxVal, dyVal); } + template __global__ + void calcMagnitudeKernel(const SrcTexObject src, PtrStepi dx, PtrStepi dy, PtrStepSzf mag, const Norm norm) + { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (y >= mag.rows || x >= mag.cols) + return; + + int dxVal = (src(y - 1, x + 1) + 2 * src(y, x + 1) + src(y + 1, x + 1)) - (src(y - 1, x - 1) + 2 * src(y, x - 1) + src(y + 1, x - 1)); + int dyVal = (src(y + 1, x - 1) + 2 * src(y + 1, x) + src(y + 1, x + 1)) - (src(y - 1, x - 1) + 2 * src(y - 1, x) + src(y - 1, x + 1)); + + dx(y, x) = dxVal; + dy(y, x) = dyVal; + + mag(y, x) = norm(dxVal, dyVal); + } + void calcMagnitude(PtrStepSzb srcWhole, int xoff, int yoff, PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, bool L2Grad, cudaStream_t stream) { const dim3 block(16, 16); const dim3 grid(divUp(mag.cols, block.x), divUp(mag.rows, block.y)); - bindTexture(&tex_src, srcWhole); - SrcTex src(xoff, yoff); + bool cc30 = deviceSupports(FEATURE_SET_COMPUTE_30); - if (L2Grad) + if (cc30) { - L2 norm; - calcMagnitudeKernel<<>>(src, dx, dy, mag, norm); + cudaResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = cudaResourceTypePitch2D; + resDesc.res.pitch2D.devPtr = srcWhole.ptr(); + resDesc.res.pitch2D.height = srcWhole.rows; + resDesc.res.pitch2D.width = srcWhole.cols; + resDesc.res.pitch2D.pitchInBytes = srcWhole.step; + resDesc.res.pitch2D.desc = cudaCreateChannelDesc(); + + cudaTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = cudaAddressModeClamp; + texDesc.addressMode[1] = cudaAddressModeClamp; + texDesc.addressMode[2] = cudaAddressModeClamp; + + cudaTextureObject_t tex = 0; + cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL); + + SrcTexObject src(xoff, yoff, tex); + + if (L2Grad) + { + L2 norm; + calcMagnitudeKernel<<>>(src, dx, dy, mag, norm); + } + else + { + L1 norm; + calcMagnitudeKernel<<>>(src, dx, dy, mag, norm); + } + + cudaSafeCall( cudaGetLastError() ); + + if (stream == NULL) + cudaSafeCall( cudaDeviceSynchronize() ); + else + cudaSafeCall( cudaStreamSynchronize(stream) ); + + cudaSafeCall( cudaDestroyTextureObject(tex) ); } else { - L1 norm; - calcMagnitudeKernel<<>>(src, dx, dy, mag, norm); + bindTexture(&tex_src, srcWhole); + SrcTex src(xoff, yoff); + + if (L2Grad) + { + L2 norm; + calcMagnitudeKernel<<>>(src, dx, dy, mag, norm); + } + else + { + L1 norm; + calcMagnitudeKernel<<>>(src, dx, dy, mag, norm); + } + + cudaSafeCall( cudaGetLastError() ); + + if (stream == NULL) + cudaSafeCall( cudaDeviceSynchronize() ); } - - cudaSafeCall( cudaGetLastError() ); - - if (stream == NULL) - cudaSafeCall( cudaDeviceSynchronize() ); } void calcMagnitude(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, bool L2Grad, cudaStream_t stream) @@ -165,7 +245,6 @@ namespace canny namespace canny { texture tex_mag(false, cudaFilterModePoint, cudaAddressModeClamp); - __global__ void calcMapKernel(const PtrStepSzi dx, const PtrStepi dy, PtrStepi map, const float low_thresh, const float high_thresh) { const int CANNY_SHIFT = 15; @@ -218,18 +297,103 @@ namespace canny map(y, x) = edge_type; } + __global__ void calcMapKernel(const PtrStepSzi dx, const PtrStepi dy, PtrStepi map, const float low_thresh, const float high_thresh, cudaTextureObject_t tex_mag) + { + const int CANNY_SHIFT = 15; + const int TG22 = (int)(0.4142135623730950488016887242097*(1<= dx.cols - 1 || y == 0 || y >= dx.rows - 1) + return; + + int dxVal = dx(y, x); + int dyVal = dy(y, x); + + const int s = (dxVal ^ dyVal) < 0 ? -1 : 1; + const float m = tex2D(tex_mag, x, y); + + dxVal = ::abs(dxVal); + dyVal = ::abs(dyVal); + + // 0 - the pixel can not belong to an edge + // 1 - the pixel might belong to an edge + // 2 - the pixel does belong to an edge + int edge_type = 0; + + if (m > low_thresh) + { + const int tg22x = dxVal * TG22; + const int tg67x = tg22x + ((dxVal + dxVal) << CANNY_SHIFT); + + dyVal <<= CANNY_SHIFT; + + if (dyVal < tg22x) + { + if (m > tex2D(tex_mag, x - 1, y) && m >= tex2D(tex_mag, x + 1, y)) + edge_type = 1 + (int)(m > high_thresh); + } + else if(dyVal > tg67x) + { + if (m > tex2D(tex_mag, x, y - 1) && m >= tex2D(tex_mag, x, y + 1)) + edge_type = 1 + (int)(m > high_thresh); + } + else + { + if (m > tex2D(tex_mag, x - s, y - 1) && m >= tex2D(tex_mag, x + s, y + 1)) + edge_type = 1 + (int)(m > high_thresh); + } + } + + map(y, x) = edge_type; + } + void calcMap(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, PtrStepSzi map, float low_thresh, float high_thresh, cudaStream_t stream) { const dim3 block(16, 16); const dim3 grid(divUp(dx.cols, block.x), divUp(dx.rows, block.y)); - bindTexture(&tex_mag, mag); + if (deviceSupports(FEATURE_SET_COMPUTE_30)) + { + // Use the texture object + cudaResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = cudaResourceTypePitch2D; + resDesc.res.pitch2D.devPtr = mag.ptr(); + resDesc.res.pitch2D.height = mag.rows; + resDesc.res.pitch2D.width = mag.cols; + resDesc.res.pitch2D.pitchInBytes = mag.step; + resDesc.res.pitch2D.desc = cudaCreateChannelDesc(); - calcMapKernel<<>>(dx, dy, map, low_thresh, high_thresh); - cudaSafeCall( cudaGetLastError() ); + cudaTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = cudaAddressModeClamp; + texDesc.addressMode[1] = cudaAddressModeClamp; + texDesc.addressMode[2] = cudaAddressModeClamp; - if (stream == NULL) - cudaSafeCall( cudaDeviceSynchronize() ); + cudaTextureObject_t tex=0; + cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL); + calcMapKernel<<>>(dx, dy, map, low_thresh, high_thresh, tex); + cudaSafeCall( cudaGetLastError() ); + + if (stream == NULL) + cudaSafeCall( cudaDeviceSynchronize() ); + else + cudaSafeCall( cudaStreamSynchronize(stream) ); + + cudaSafeCall( cudaDestroyTextureObject(tex) ); + } + else + { + // Use the texture reference + bindTexture(&tex_mag, mag); + calcMapKernel<<>>(dx, dy, map, low_thresh, high_thresh); + cudaSafeCall( cudaGetLastError() ); + + if (stream == NULL) + cudaSafeCall( cudaDeviceSynchronize() ); + } } } @@ -237,14 +401,12 @@ namespace canny namespace canny { - __device__ int counter = 0; - __device__ __forceinline__ bool checkIdx(int y, int x, int rows, int cols) { return (y >= 0) && (y < rows) && (x >= 0) && (x < cols); } - __global__ void edgesHysteresisLocalKernel(PtrStepSzi map, short2* st) + __global__ void edgesHysteresisLocalKernel(PtrStepSzi map, short2* st, int* d_counter) { __shared__ volatile int smem[18][18]; @@ -325,22 +487,19 @@ namespace canny if (n > 0) { - const int ind = ::atomicAdd(&counter, 1); + const int ind = ::atomicAdd(d_counter, 1); st[ind] = make_short2(x, y); } } - void edgesHysteresisLocal(PtrStepSzi map, short2* st1, cudaStream_t stream) + void edgesHysteresisLocal(PtrStepSzi map, short2* st1, int* d_counter, cudaStream_t stream) { - void* counter_ptr; - cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) ); - - cudaSafeCall( cudaMemsetAsync(counter_ptr, 0, sizeof(int), stream) ); + cudaSafeCall( cudaMemsetAsync(d_counter, 0, sizeof(int), stream) ); const dim3 block(16, 16); const dim3 grid(divUp(map.cols, block.x), divUp(map.rows, block.y)); - edgesHysteresisLocalKernel<<>>(map, st1); + edgesHysteresisLocalKernel<<>>(map, st1, d_counter); cudaSafeCall( cudaGetLastError() ); if (stream == NULL) @@ -355,7 +514,7 @@ namespace canny __constant__ int c_dx[8] = {-1, 0, 1, -1, 1, -1, 0, 1}; __constant__ int c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1}; - __global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, short2* st1, short2* st2, const int count) + __global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, short2* st1, short2* st2, int* d_counter, const int count) { const int stack_size = 512; @@ -429,7 +588,7 @@ namespace canny { if (threadIdx.x == 0) { - s_ind = ::atomicAdd(&counter, s_counter); + s_ind = ::atomicAdd(d_counter, s_counter); if (s_ind + s_counter > map.cols * map.rows) s_counter = 0; @@ -444,29 +603,26 @@ namespace canny } } - void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2, cudaStream_t stream) + void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2, int* d_counter, cudaStream_t stream) { - void* counter_ptr; - cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) ); - int count; - cudaSafeCall( cudaMemcpyAsync(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost, stream) ); + cudaSafeCall( cudaMemcpyAsync(&count, d_counter, sizeof(int), cudaMemcpyDeviceToHost, stream) ); cudaSafeCall( cudaStreamSynchronize(stream) ); while (count > 0) { - cudaSafeCall( cudaMemsetAsync(counter_ptr, 0, sizeof(int), stream) ); + cudaSafeCall( cudaMemsetAsync(d_counter, 0, sizeof(int), stream) ); const dim3 block(128); const dim3 grid(::min(count, 65535u), divUp(count, 65535), 1); - edgesHysteresisGlobalKernel<<>>(map, st1, st2, count); + edgesHysteresisGlobalKernel<<>>(map, st1, st2, d_counter, count); cudaSafeCall( cudaGetLastError() ); if (stream == NULL) cudaSafeCall( cudaDeviceSynchronize() ); - cudaSafeCall( cudaMemcpyAsync(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost, stream) ); + cudaSafeCall( cudaMemcpyAsync(&count, d_counter, sizeof(int), cudaMemcpyDeviceToHost, stream) ); cudaSafeCall( cudaStreamSynchronize(stream) ); count = min(count, map.cols * map.rows); diff --git a/modules/cudaimgproc/test/test_canny.cpp b/modules/cudaimgproc/test/test_canny.cpp index 2b1a8d57e0..a782a87b3b 100644 --- a/modules/cudaimgproc/test/test_canny.cpp +++ b/modules/cudaimgproc/test/test_canny.cpp @@ -92,9 +92,66 @@ CUDA_TEST_P(Canny, Accuracy) EXPECT_MAT_SIMILAR(edges_gold, edges, 2e-2); } +class CannyAsyncParallelLoopBody : public cv::ParallelLoopBody +{ +public: + CannyAsyncParallelLoopBody(const cv::cuda::GpuMat& d_img_, cv::cuda::GpuMat* edges_, double low_thresh_, double high_thresh_, int apperture_size_, bool useL2gradient_) + : d_img(d_img_), edges(edges_), low_thresh(low_thresh_), high_thresh(high_thresh_), apperture_size(apperture_size_), useL2gradient(useL2gradient_) {} + ~CannyAsyncParallelLoopBody() {}; + void operator()(const cv::Range& r) const + { + for (int i = r.start; i < r.end; i++) { + cv::cuda::Stream stream; + cv::Ptr canny = cv::cuda::createCannyEdgeDetector(low_thresh, high_thresh, apperture_size, useL2gradient); + canny->detect(d_img, edges[i], stream); + stream.waitForCompletion(); + } + } +protected: + const cv::cuda::GpuMat& d_img; + cv::cuda::GpuMat* edges; + double low_thresh; + double high_thresh; + int apperture_size; + bool useL2gradient; +}; + +#define NUM_STREAMS 64 + +CUDA_TEST_P(Canny, Async) +{ + if (!supportFeature(devInfo, cv::cuda::FEATURE_SET_COMPUTE_30)) + { + throw SkipTestException("CUDA device doesn't support texture objects"); + } + else + { + const cv::Mat img = readImage("stereobm/aloe-L.png", cv::IMREAD_GRAYSCALE); + ASSERT_FALSE(img.empty()); + + const cv::cuda::GpuMat d_img_roi = loadMat(img, useRoi); + + double low_thresh = 50.0; + double high_thresh = 100.0; + + // Synchronous call + cv::Ptr canny = cv::cuda::createCannyEdgeDetector(low_thresh, high_thresh, apperture_size, useL2gradient); + cv::cuda::GpuMat edges_gold; + canny->detect(d_img_roi, edges_gold); + + // Asynchronous call + cv::cuda::GpuMat edges[NUM_STREAMS]; + cv::parallel_for_(cv::Range(0, NUM_STREAMS), CannyAsyncParallelLoopBody(d_img_roi, edges, low_thresh, high_thresh, apperture_size, useL2gradient)); + + // Compare the results of synchronous call and asynchronous call + for (int i = 0; i < NUM_STREAMS; i++) + EXPECT_MAT_NEAR(edges_gold, edges[i], 0.0); + } + } + INSTANTIATE_TEST_CASE_P(CUDA_ImgProc, Canny, testing::Combine( ALL_DEVICES, - testing::Values(AppertureSize(3), AppertureSize(5)), + testing::Values(AppertureSize(3), AppertureSize(5), AppertureSize(7)), testing::Values(L2gradient(false), L2gradient(true)), WHOLE_SUBMAT));