arm_compute v19.08

commit: 975dfe175e3d7c62c27598b1c0e8e77ed90df463 [log] [tgz]
author: Jenkins <bsgcomp@arm.com> Mon Sep 02 11:47:54 2019 +0100
committer: Jenkins <bsgcomp@arm.com> Mon Sep 02 11:47:54 2019 +0100
tree: 3a7a8bd660c08b4b2ba68ed6a260ac2dd440c3ad
parent: 4ba87dbdc3b22220eba4a792c1f5c87e7a88c7af [diff] [blame]
diff --git a/documentation/minmaxloc_8cl.xhtml b/documentation/minmaxloc_8cl.xhtml
index cc31473..f1c8ce6 100644
--- a/documentation/minmaxloc_8cl.xhtml
+++ b/documentation/minmaxloc_8cl.xhtml

@@ -40,7 +40,7 @@
   <img alt="Compute Library" src="https://raw.githubusercontent.com/ARM-software/ComputeLibrary/gh-pages/ACL_logo.png" style="max-width: 100%;margin-top: 15px;margin-left: 10px"/>
   <td style="padding-left: 0.5em;">
    <div id="projectname">
-   &#160;<span id="projectnumber">19.05</span>
+   &#160;<span id="projectnumber">19.08</span>
    </div>
   </td>
  </tr>
@@ -281,21 +281,21 @@
 </dl>
 
 <p class="definition">Definition at line <a class="el" href="minmaxloc_8cl_source.xhtml#l00065">65</a> of file <a class="el" href="minmaxloc_8cl_source.xhtml">minmaxloc.cl</a>.</p>
-<div class="fragment"><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;{</div><div class="line"><a name="l00070"></a><span class="lineno">   70</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00071"></a><span class="lineno">   71</span>&#160;</div><div class="line"><a name="l00072"></a><span class="lineno">   72</span>&#160;    <span class="comment">// Initialize local minimum and local maximum</span></div><div class="line"><a name="l00073"></a><span class="lineno">   73</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00074"></a><span class="lineno">   74</span>&#160;    local_min = <a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>;</div><div class="line"><a name="l00075"></a><span class="lineno">   75</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00076"></a><span class="lineno">   76</span>&#160;    local_max = <a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>;</div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;</div><div class="line"><a name="l00078"></a><span class="lineno">   78</span>&#160;    <span class="comment">// Calculate min/max of row</span></div><div class="line"><a name="l00079"></a><span class="lineno">   79</span>&#160;    <span class="keywordtype">int</span> i = 0;</div><div class="line"><a name="l00080"></a><span class="lineno">   80</span>&#160;    <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">for</a>(; i + 16 &lt;= width; i += 16)</div><div class="line"><a name="l00081"></a><span class="lineno">   81</span>&#160;    {</div><div class="line"><a name="l00082"></a><span class="lineno">   82</span>&#160;        <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00083"></a><span class="lineno">   83</span>&#160;        data      = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i, 0));</div><div class="line"><a name="l00084"></a><span class="lineno">   84</span>&#160;        local_min = min(data, local_min);</div><div class="line"><a name="l00085"></a><span class="lineno">   85</span>&#160;        local_max = max(data, local_max);</div><div class="line"><a name="l00086"></a><span class="lineno">   86</span>&#160;    }</div><div class="line"><a name="l00087"></a><span class="lineno">   87</span>&#160;</div><div class="line"><a name="l00088"></a><span class="lineno">   88</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00090"></a><span class="lineno">   90</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00091"></a><span class="lineno">   91</span>&#160;    data = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i, 0));</div><div class="line"><a name="l00092"></a><span class="lineno">   92</span>&#160;<span class="preprocessor">#ifdef IS_DATA_TYPE_FLOAT</span></div><div class="line"><a name="l00093"></a><span class="lineno">   93</span>&#160;    int16 valid_indices = (i + <a class="code" href="minmaxloc_8cl.xhtml#a5a67e1e00d711339bef1e10f0984227c">idx16</a>) &lt; width;</div><div class="line"><a name="l00094"></a><span class="lineno">   94</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00095"></a><span class="lineno">   95</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00096"></a><span class="lineno">   96</span>&#160;    valid_indices = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((i + <a class="code" href="minmaxloc_8cl.xhtml#a5a67e1e00d711339bef1e10f0984227c">idx16</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16));</div><div class="line"><a name="l00097"></a><span class="lineno">   97</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00098"></a><span class="lineno">   98</span>&#160;    local_max = max(local_max, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(<a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>, data, valid_indices));</div><div class="line"><a name="l00099"></a><span class="lineno">   99</span>&#160;    local_min = min(local_min, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(<a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>, data, valid_indices));</div><div class="line"><a name="l00100"></a><span class="lineno">  100</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_16 */</span><span class="preprocessor"></span></div><div class="line"><a name="l00101"></a><span class="lineno">  101</span>&#160;</div><div class="line"><a name="l00102"></a><span class="lineno">  102</span>&#160;    <span class="comment">// Perform min/max reduction</span></div><div class="line"><a name="l00103"></a><span class="lineno">  103</span>&#160;    local_min.s01234567 = min(local_min.s01234567, local_min.s89ABCDEF);</div><div class="line"><a name="l00104"></a><span class="lineno">  104</span>&#160;    local_max.s01234567 = max(local_max.s01234567, local_max.s89ABCDEF);</div><div class="line"><a name="l00105"></a><span class="lineno">  105</span>&#160;</div><div class="line"><a name="l00106"></a><span class="lineno">  106</span>&#160;    local_min.s0123 = min(local_min.s0123, local_min.s4567);</div><div class="line"><a name="l00107"></a><span class="lineno">  107</span>&#160;    local_max.s0123 = max(local_max.s0123, local_max.s4567);</div><div class="line"><a name="l00108"></a><span class="lineno">  108</span>&#160;</div><div class="line"><a name="l00109"></a><span class="lineno">  109</span>&#160;    local_min.s01 = min(local_min.s01, local_min.s23);</div><div class="line"><a name="l00110"></a><span class="lineno">  110</span>&#160;    local_max.s01 = max(local_max.s01, local_max.s23);</div><div class="line"><a name="l00111"></a><span class="lineno">  111</span>&#160;</div><div class="line"><a name="l00112"></a><span class="lineno">  112</span>&#160;    local_min.s0 = min(local_min.s0, local_min.s1);</div><div class="line"><a name="l00113"></a><span class="lineno">  113</span>&#160;    local_max.s0 = max(local_max.s0, local_max.s1);</div><div class="line"><a name="l00114"></a><span class="lineno">  114</span>&#160;</div><div class="line"><a name="l00115"></a><span class="lineno">  115</span>&#160;    <span class="comment">// Update global min/max</span></div><div class="line"><a name="l00116"></a><span class="lineno">  116</span>&#160;<span class="preprocessor">#ifdef IS_DATA_TYPE_FLOAT</span></div><div class="line"><a name="l00117"></a><span class="lineno">  117</span>&#160;    atomic_min(&amp;min_max[0], <a class="code" href="minmaxloc_8cl.xhtml#a061903419fb2aff60da86e674c76b96e">FloatFlip</a>(local_min.s0));</div><div class="line"><a name="l00118"></a><span class="lineno">  118</span>&#160;    atomic_max(&amp;min_max[1], <a class="code" href="minmaxloc_8cl.xhtml#a061903419fb2aff60da86e674c76b96e">FloatFlip</a>(local_max.s0));</div><div class="line"><a name="l00119"></a><span class="lineno">  119</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00120"></a><span class="lineno">  120</span>&#160;    atomic_min(&amp;min_max[0], local_min.s0);</div><div class="line"><a name="l00121"></a><span class="lineno">  121</span>&#160;    atomic_max(&amp;min_max[1], local_max.s0);</div><div class="line"><a name="l00122"></a><span class="lineno">  122</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00123"></a><span class="lineno">  123</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00312">helpers.h:312</a></div></div>
-<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00063">helpers.h:63</a></div></div>
-<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00113">helpers.h:113</a></div></div>
+<div class="fragment"><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;{</div><div class="line"><a name="l00070"></a><span class="lineno">   70</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00071"></a><span class="lineno">   71</span>&#160;</div><div class="line"><a name="l00072"></a><span class="lineno">   72</span>&#160;    <span class="comment">// Initialize local minimum and local maximum</span></div><div class="line"><a name="l00073"></a><span class="lineno">   73</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00074"></a><span class="lineno">   74</span>&#160;    local_min = <a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>;</div><div class="line"><a name="l00075"></a><span class="lineno">   75</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00076"></a><span class="lineno">   76</span>&#160;    local_max = <a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>;</div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;</div><div class="line"><a name="l00078"></a><span class="lineno">   78</span>&#160;    <span class="comment">// Calculate min/max of row</span></div><div class="line"><a name="l00079"></a><span class="lineno">   79</span>&#160;    <span class="keywordtype">int</span> i = 0;</div><div class="line"><a name="l00080"></a><span class="lineno">   80</span>&#160;    <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">for</a>(; i + 16 &lt;= width; i += 16)</div><div class="line"><a name="l00081"></a><span class="lineno">   81</span>&#160;    {</div><div class="line"><a name="l00082"></a><span class="lineno">   82</span>&#160;        <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00083"></a><span class="lineno">   83</span>&#160;        data      = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i, 0));</div><div class="line"><a name="l00084"></a><span class="lineno">   84</span>&#160;        local_min = min(data, local_min);</div><div class="line"><a name="l00085"></a><span class="lineno">   85</span>&#160;        local_max = max(data, local_max);</div><div class="line"><a name="l00086"></a><span class="lineno">   86</span>&#160;    }</div><div class="line"><a name="l00087"></a><span class="lineno">   87</span>&#160;</div><div class="line"><a name="l00088"></a><span class="lineno">   88</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00090"></a><span class="lineno">   90</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00091"></a><span class="lineno">   91</span>&#160;    data = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i, 0));</div><div class="line"><a name="l00092"></a><span class="lineno">   92</span>&#160;<span class="preprocessor">#ifdef IS_DATA_TYPE_FLOAT</span></div><div class="line"><a name="l00093"></a><span class="lineno">   93</span>&#160;    int16 valid_indices = (i + <a class="code" href="minmaxloc_8cl.xhtml#a5a67e1e00d711339bef1e10f0984227c">idx16</a>) &lt; width;</div><div class="line"><a name="l00094"></a><span class="lineno">   94</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00095"></a><span class="lineno">   95</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00096"></a><span class="lineno">   96</span>&#160;    valid_indices = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((i + <a class="code" href="minmaxloc_8cl.xhtml#a5a67e1e00d711339bef1e10f0984227c">idx16</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16));</div><div class="line"><a name="l00097"></a><span class="lineno">   97</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00098"></a><span class="lineno">   98</span>&#160;    local_max = max(local_max, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(<a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>, data, valid_indices));</div><div class="line"><a name="l00099"></a><span class="lineno">   99</span>&#160;    local_min = min(local_min, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(<a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>, data, valid_indices));</div><div class="line"><a name="l00100"></a><span class="lineno">  100</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_16 */</span><span class="preprocessor"></span></div><div class="line"><a name="l00101"></a><span class="lineno">  101</span>&#160;</div><div class="line"><a name="l00102"></a><span class="lineno">  102</span>&#160;    <span class="comment">// Perform min/max reduction</span></div><div class="line"><a name="l00103"></a><span class="lineno">  103</span>&#160;    local_min.s01234567 = min(local_min.s01234567, local_min.s89ABCDEF);</div><div class="line"><a name="l00104"></a><span class="lineno">  104</span>&#160;    local_max.s01234567 = max(local_max.s01234567, local_max.s89ABCDEF);</div><div class="line"><a name="l00105"></a><span class="lineno">  105</span>&#160;</div><div class="line"><a name="l00106"></a><span class="lineno">  106</span>&#160;    local_min.s0123 = min(local_min.s0123, local_min.s4567);</div><div class="line"><a name="l00107"></a><span class="lineno">  107</span>&#160;    local_max.s0123 = max(local_max.s0123, local_max.s4567);</div><div class="line"><a name="l00108"></a><span class="lineno">  108</span>&#160;</div><div class="line"><a name="l00109"></a><span class="lineno">  109</span>&#160;    local_min.s01 = min(local_min.s01, local_min.s23);</div><div class="line"><a name="l00110"></a><span class="lineno">  110</span>&#160;    local_max.s01 = max(local_max.s01, local_max.s23);</div><div class="line"><a name="l00111"></a><span class="lineno">  111</span>&#160;</div><div class="line"><a name="l00112"></a><span class="lineno">  112</span>&#160;    local_min.s0 = min(local_min.s0, local_min.s1);</div><div class="line"><a name="l00113"></a><span class="lineno">  113</span>&#160;    local_max.s0 = max(local_max.s0, local_max.s1);</div><div class="line"><a name="l00114"></a><span class="lineno">  114</span>&#160;</div><div class="line"><a name="l00115"></a><span class="lineno">  115</span>&#160;    <span class="comment">// Update global min/max</span></div><div class="line"><a name="l00116"></a><span class="lineno">  116</span>&#160;<span class="preprocessor">#ifdef IS_DATA_TYPE_FLOAT</span></div><div class="line"><a name="l00117"></a><span class="lineno">  117</span>&#160;    atomic_min(&amp;min_max[0], <a class="code" href="minmaxloc_8cl.xhtml#a061903419fb2aff60da86e674c76b96e">FloatFlip</a>(local_min.s0));</div><div class="line"><a name="l00118"></a><span class="lineno">  118</span>&#160;    atomic_max(&amp;min_max[1], <a class="code" href="minmaxloc_8cl.xhtml#a061903419fb2aff60da86e674c76b96e">FloatFlip</a>(local_max.s0));</div><div class="line"><a name="l00119"></a><span class="lineno">  119</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00120"></a><span class="lineno">  120</span>&#160;    atomic_min(&amp;min_max[0], local_min.s0);</div><div class="line"><a name="l00121"></a><span class="lineno">  121</span>&#160;    atomic_max(&amp;min_max[1], local_max.s0);</div><div class="line"><a name="l00122"></a><span class="lineno">  122</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00123"></a><span class="lineno">  123</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00328">helpers.h:328</a></div></div>
+<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00079">helpers.h:79</a></div></div>
+<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00129">helpers.h:129</a></div></div>
 <div class="ttc" id="minmaxloc_8cl_xhtml_a5a67e1e00d711339bef1e10f0984227c"><div class="ttname"><a href="minmaxloc_8cl.xhtml#a5a67e1e00d711339bef1e10f0984227c">idx16</a></div><div class="ttdeci">__constant int16 idx16</div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00048">minmaxloc.cl:48</a></div></div>
 <div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a0886942393a3ba0dfefaa7516b159784"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">arm_compute::test::validation::for</a></div><div class="ttdeci">for(size_t k=0;k&lt; _target.size();++k)</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_unstack_8cpp_source.xhtml#l00091">Unstack.cpp:91</a></div></div>
-<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00151">helpers.h:151</a></div></div>
+<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00167">helpers.h:167</a></div></div>
 <div class="ttc" id="minmaxloc_8cl_xhtml_a538b4b63f40e7b12891774e03a4f0dec"><div class="ttname"><a href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a></div><div class="ttdeci">__constant DATA_TYPE16 type_min</div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00046">minmaxloc.cl:46</a></div></div>
 <div class="ttc" id="minmaxloc_8cl_xhtml_a061903419fb2aff60da86e674c76b96e"><div class="ttname"><a href="minmaxloc_8cl.xhtml#a061903419fb2aff60da86e674c76b96e">FloatFlip</a></div><div class="ttdeci">int FloatFlip(float val)</div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00035">minmaxloc.cl:35</a></div></div>
 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_af77145fbdc6b0c8931148f5597d9de53"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">arm_compute::test::validation::select</a></div><div class="ttdeci">CLSelect select</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_select_8cpp_source.xhtml#l00164">Select.cpp:164</a></div></div>
 <div class="ttc" id="minmaxloc_8cl_xhtml_a4464d6f922ea17b4a9ca6a2cec7ddb75"><div class="ttname"><a href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a></div><div class="ttdeci">__constant DATA_TYPE16 type_max</div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00047">minmaxloc.cl:47</a></div></div>
-<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
-<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00057">helpers.h:57</a></div></div>
+<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
+<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00073">helpers.h:73</a></div></div>
 </div><!-- fragment -->
-<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00063">CONVERT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00113">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00035">FloatFlip()</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00048">idx16</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00312">offset()</a>, <a class="el" href="_c_l_2_select_8cpp_source.xhtml#l00164">arm_compute::test::validation::select</a>, <a class="el" href="_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00047">type_max</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00046">type_min</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00057">VEC_DATA_TYPE</a>.</p>
+<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00079">CONVERT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00129">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00035">FloatFlip()</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00048">idx16</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00328">offset()</a>, <a class="el" href="_c_l_2_select_8cpp_source.xhtml#l00164">arm_compute::test::validation::select</a>, <a class="el" href="_c_l_2_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00047">type_max</a>, <a class="el" href="minmaxloc_8cl_source.xhtml#l00046">type_min</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00073">VEC_DATA_TYPE</a>.</p>
 
 </div>
 </div>
@@ -408,12 +408,12 @@
 </dl>
 
 <p class="definition">Definition at line <a class="el" href="minmaxloc_8cl_source.xhtml#l00143">143</a> of file <a class="el" href="minmaxloc_8cl_source.xhtml">minmaxloc.cl</a>.</p>
-<div class="fragment"><div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;{</div><div class="line"><a name="l00157"></a><span class="lineno">  157</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00158"></a><span class="lineno">  158</span>&#160;</div><div class="line"><a name="l00159"></a><span class="lineno">  159</span>&#160;<span class="preprocessor">#ifdef IS_DATA_TYPE_FLOAT</span></div><div class="line"><a name="l00160"></a><span class="lineno">  160</span>&#160;    __global <span class="keywordtype">float</span> *min_max_ptr = (__global <span class="keywordtype">float</span> *)min_max;</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;    <span class="keywordtype">float</span>           min_value   = min_max_ptr[0];</div><div class="line"><a name="l00162"></a><span class="lineno">  162</span>&#160;    <span class="keywordtype">float</span>           max_value   = min_max_ptr[1];</div><div class="line"><a name="l00163"></a><span class="lineno">  163</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00164"></a><span class="lineno">  164</span>&#160;    <span class="keywordtype">int</span> min_value = min_max[0];</div><div class="line"><a name="l00165"></a><span class="lineno">  165</span>&#160;    <span class="keywordtype">int</span> max_value = min_max[1];</div><div class="line"><a name="l00166"></a><span class="lineno">  166</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00167"></a><span class="lineno">  167</span>&#160;</div><div class="line"><a name="l00168"></a><span class="lineno">  168</span>&#160;    <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> value = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>.ptr);</div><div class="line"><a name="l00169"></a><span class="lineno">  169</span>&#160;<span class="preprocessor">#ifdef COUNT_MIN_MAX</span></div><div class="line"><a name="l00170"></a><span class="lineno">  170</span>&#160;    <span class="keywordflow">if</span>(value == min_value)</div><div class="line"><a name="l00171"></a><span class="lineno">  171</span>&#160;    {</div><div class="line"><a name="l00172"></a><span class="lineno">  172</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[0]);</div><div class="line"><a name="l00173"></a><span class="lineno">  173</span>&#160;<span class="preprocessor">#ifdef LOCATE_MIN</span></div><div class="line"><a name="l00174"></a><span class="lineno">  174</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_min_loc_count)</div><div class="line"><a name="l00175"></a><span class="lineno">  175</span>&#160;        {</div><div class="line"><a name="l00176"></a><span class="lineno">  176</span>&#160;            min_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00177"></a><span class="lineno">  177</span>&#160;            min_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00178"></a><span class="lineno">  178</span>&#160;        }</div><div class="line"><a name="l00179"></a><span class="lineno">  179</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOCATE_MIN */</span><span class="preprocessor"></span></div><div class="line"><a name="l00180"></a><span class="lineno">  180</span>&#160;    }</div><div class="line"><a name="l00181"></a><span class="lineno">  181</span>&#160;    <span class="keywordflow">if</span>(value == max_value)</div><div class="line"><a name="l00182"></a><span class="lineno">  182</span>&#160;    {</div><div class="line"><a name="l00183"></a><span class="lineno">  183</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[1]);</div><div class="line"><a name="l00184"></a><span class="lineno">  184</span>&#160;<span class="preprocessor">#ifdef LOCATE_MAX</span></div><div class="line"><a name="l00185"></a><span class="lineno">  185</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_max_loc_count)</div><div class="line"><a name="l00186"></a><span class="lineno">  186</span>&#160;        {</div><div class="line"><a name="l00187"></a><span class="lineno">  187</span>&#160;            max_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00188"></a><span class="lineno">  188</span>&#160;            max_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00189"></a><span class="lineno">  189</span>&#160;        }</div><div class="line"><a name="l00190"></a><span class="lineno">  190</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOCATE_MAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00191"></a><span class="lineno">  191</span>&#160;    }</div><div class="line"><a name="l00192"></a><span class="lineno">  192</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* COUNT_MIN_MAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00193"></a><span class="lineno">  193</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00113">helpers.h:113</a></div></div>
+<div class="fragment"><div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;{</div><div class="line"><a name="l00157"></a><span class="lineno">  157</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00158"></a><span class="lineno">  158</span>&#160;</div><div class="line"><a name="l00159"></a><span class="lineno">  159</span>&#160;<span class="preprocessor">#ifdef IS_DATA_TYPE_FLOAT</span></div><div class="line"><a name="l00160"></a><span class="lineno">  160</span>&#160;    __global <span class="keywordtype">float</span> *min_max_ptr = (__global <span class="keywordtype">float</span> *)min_max;</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;    <span class="keywordtype">float</span>           min_value   = min_max_ptr[0];</div><div class="line"><a name="l00162"></a><span class="lineno">  162</span>&#160;    <span class="keywordtype">float</span>           max_value   = min_max_ptr[1];</div><div class="line"><a name="l00163"></a><span class="lineno">  163</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00164"></a><span class="lineno">  164</span>&#160;    <span class="keywordtype">int</span> min_value = min_max[0];</div><div class="line"><a name="l00165"></a><span class="lineno">  165</span>&#160;    <span class="keywordtype">int</span> max_value = min_max[1];</div><div class="line"><a name="l00166"></a><span class="lineno">  166</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* IS_DATA_TYPE_FLOAT */</span><span class="preprocessor"></span></div><div class="line"><a name="l00167"></a><span class="lineno">  167</span>&#160;</div><div class="line"><a name="l00168"></a><span class="lineno">  168</span>&#160;    <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> value = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>.ptr);</div><div class="line"><a name="l00169"></a><span class="lineno">  169</span>&#160;<span class="preprocessor">#ifdef COUNT_MIN_MAX</span></div><div class="line"><a name="l00170"></a><span class="lineno">  170</span>&#160;    <span class="keywordflow">if</span>(value == min_value)</div><div class="line"><a name="l00171"></a><span class="lineno">  171</span>&#160;    {</div><div class="line"><a name="l00172"></a><span class="lineno">  172</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[0]);</div><div class="line"><a name="l00173"></a><span class="lineno">  173</span>&#160;<span class="preprocessor">#ifdef LOCATE_MIN</span></div><div class="line"><a name="l00174"></a><span class="lineno">  174</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_min_loc_count)</div><div class="line"><a name="l00175"></a><span class="lineno">  175</span>&#160;        {</div><div class="line"><a name="l00176"></a><span class="lineno">  176</span>&#160;            min_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00177"></a><span class="lineno">  177</span>&#160;            min_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00178"></a><span class="lineno">  178</span>&#160;        }</div><div class="line"><a name="l00179"></a><span class="lineno">  179</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOCATE_MIN */</span><span class="preprocessor"></span></div><div class="line"><a name="l00180"></a><span class="lineno">  180</span>&#160;    }</div><div class="line"><a name="l00181"></a><span class="lineno">  181</span>&#160;    <span class="keywordflow">if</span>(value == max_value)</div><div class="line"><a name="l00182"></a><span class="lineno">  182</span>&#160;    {</div><div class="line"><a name="l00183"></a><span class="lineno">  183</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[1]);</div><div class="line"><a name="l00184"></a><span class="lineno">  184</span>&#160;<span class="preprocessor">#ifdef LOCATE_MAX</span></div><div class="line"><a name="l00185"></a><span class="lineno">  185</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_max_loc_count)</div><div class="line"><a name="l00186"></a><span class="lineno">  186</span>&#160;        {</div><div class="line"><a name="l00187"></a><span class="lineno">  187</span>&#160;            max_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00188"></a><span class="lineno">  188</span>&#160;            max_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00189"></a><span class="lineno">  189</span>&#160;        }</div><div class="line"><a name="l00190"></a><span class="lineno">  190</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOCATE_MAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00191"></a><span class="lineno">  191</span>&#160;    }</div><div class="line"><a name="l00192"></a><span class="lineno">  192</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* COUNT_MIN_MAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00193"></a><span class="lineno">  193</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00129">helpers.h:129</a></div></div>
 <div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
-<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00151">helpers.h:151</a></div></div>
-<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
+<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00167">helpers.h:167</a></div></div>
+<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
 </div><!-- fragment -->
-<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00113">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, and <a class="el" href="_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>.</p>
+<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00129">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, and <a class="el" href="_c_l_2_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>.</p>
 
 </div>
 </div>
@@ -478,7 +478,7 @@
 <div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
   <ul>
     <li class="navelem"><a class="el" href="dir_68267d1309a1af8e8297ef4c3efbcdba.xhtml">src</a></li><li class="navelem"><a class="el" href="dir_aebb8dcc11953d78e620bbef0b9e2183.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_8c278f79c760e5c5fbd911f9870614c1.xhtml">CL</a></li><li class="navelem"><a class="el" href="dir_25885286e9dad4fa105b7b25a8031bbf.xhtml">cl_kernels</a></li><li class="navelem"><a class="el" href="minmaxloc_8cl.xhtml">minmaxloc.cl</a></li>
-    <li class="footer">Generated on Thu May 23 2019 17:11:26 for Compute Library by
+    <li class="footer">Generated on Mon Sep 2 2019 11:47:29 for Compute Library by
     <a href="http://www.doxygen.org/index.html">
     <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.15 </li>
   </ul>
commit	975dfe175e3d7c62c27598b1c0e8e77ed90df463	[log] [tgz]
author	Jenkins <bsgcomp@arm.com>	Mon Sep 02 11:47:54 2019 +0100
committer	Jenkins <bsgcomp@arm.com>	Mon Sep 02 11:47:54 2019 +0100
tree	3a7a8bd660c08b4b2ba68ed6a260ac2dd440c3ad
parent	4ba87dbdc3b22220eba4a792c1f5c87e7a88c7af [diff] [blame]