arm_compute v17.06
diff --git a/documentation/minmaxloc_8cl_source.xhtml b/documentation/minmaxloc_8cl_source.xhtml
index bf06697..7998630 100644
--- a/documentation/minmaxloc_8cl_source.xhtml
+++ b/documentation/minmaxloc_8cl_source.xhtml
@@ -6,7 +6,7 @@
 <meta http-equiv="X-UA-Compatible" content="IE=9"/>
 <meta name="generator" content="Doxygen 1.8.11"/>
 <meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
-<title>ARM Compute Library: src/core/CL/cl_kernels/minmaxloc.cl Source File</title>
+<title>Compute Library: src/core/CL/cl_kernels/minmaxloc.cl Source File</title>
 <link href="tabs.css" rel="stylesheet" type="text/css"/>
 <script type="text/javascript" src="jquery.js"></script>
 <script type="text/javascript" src="dynsections.js"></script>
@@ -39,8 +39,8 @@
  <tbody>
  <tr style="height: 56px;">
   <td style="padding-left: 0.5em;">
-   <div id="projectname">ARM Compute Library
-   &#160;<span id="projectnumber">17.05</span>
+   <div id="projectname">Compute Library
+   &#160;<span id="projectnumber">v17.06</span>
    </div>
   </td>
  </tr>
@@ -55,6 +55,7 @@
   <div id="navrow1" class="tabs">
     <ul class="tablist">
       <li><a href="index.xhtml"><span>Main&#160;Page</span></a></li>
+      <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
       <li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
       <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
       <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
@@ -116,7 +117,8 @@
 <div class="title">minmaxloc.cl</div>  </div>
 </div><!--header-->
 <div class="contents">
-<a href="minmaxloc_8cl.xhtml">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno">    1</span>&#160;<span class="comment">/*</span></div><div class="line"><a name="l00002"></a><span class="lineno">    2</span>&#160;<span class="comment"> * Copyright (c) 2017 ARM Limited.</span></div><div class="line"><a name="l00003"></a><span class="lineno">    3</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno">    4</span>&#160;<span class="comment"> * SPDX-License-Identifier: MIT</span></div><div class="line"><a name="l00005"></a><span class="lineno">    5</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00006"></a><span class="lineno">    6</span>&#160;<span class="comment"> * Permission is hereby granted, free of charge, to any person obtaining a copy</span></div><div class="line"><a name="l00007"></a><span class="lineno">    7</span>&#160;<span class="comment"> * of this software and associated documentation files (the &quot;Software&quot;), to</span></div><div class="line"><a name="l00008"></a><span class="lineno">    8</span>&#160;<span class="comment"> * deal in the Software without restriction, including without limitation the</span></div><div class="line"><a name="l00009"></a><span class="lineno">    9</span>&#160;<span class="comment"> * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or</span></div><div class="line"><a name="l00010"></a><span class="lineno">   10</span>&#160;<span class="comment"> * sell copies of the Software, and to permit persons to whom the Software is</span></div><div class="line"><a name="l00011"></a><span class="lineno">   11</span>&#160;<span class="comment"> * furnished to do so, subject to the following conditions:</span></div><div class="line"><a name="l00012"></a><span class="lineno">   12</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00013"></a><span class="lineno">   13</span>&#160;<span class="comment"> * The above copyright notice and this permission notice shall be included in all</span></div><div class="line"><a name="l00014"></a><span class="lineno">   14</span>&#160;<span class="comment"> * copies or substantial portions of the Software.</span></div><div class="line"><a name="l00015"></a><span class="lineno">   15</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00016"></a><span class="lineno">   16</span>&#160;<span class="comment"> * THE SOFTWARE IS PROVIDED &quot;AS IS&quot;, WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</span></div><div class="line"><a name="l00017"></a><span class="lineno">   17</span>&#160;<span class="comment"> * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,</span></div><div class="line"><a name="l00018"></a><span class="lineno">   18</span>&#160;<span class="comment"> * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE</span></div><div class="line"><a name="l00019"></a><span class="lineno">   19</span>&#160;<span class="comment"> * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER</span></div><div class="line"><a name="l00020"></a><span class="lineno">   20</span>&#160;<span class="comment"> * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,</span></div><div class="line"><a name="l00021"></a><span class="lineno">   21</span>&#160;<span class="comment"> * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE</span></div><div class="line"><a name="l00022"></a><span class="lineno">   22</span>&#160;<span class="comment"> * SOFTWARE.</span></div><div class="line"><a name="l00023"></a><span class="lineno">   23</span>&#160;<span class="comment"> */</span></div><div class="line"><a name="l00024"></a><span class="lineno">   24</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="helpers_8h.xhtml">helpers.h</a>&quot;</span></div><div class="line"><a name="l00025"></a><span class="lineno">   25</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="types_8h.xhtml">types.h</a>&quot;</span></div><div class="line"><a name="l00026"></a><span class="lineno">   26</span>&#160;</div><div class="line"><a name="l00027"></a><span class="lineno">   27</span>&#160;<span class="preprocessor">#ifndef DATA_TYPE_MIN</span></div><div class="line"><a name="l00028"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a17decce255d6f94a60379a61a54e3bac">   28</a></span>&#160;<span class="preprocessor">#define DATA_TYPE_MIN 0x0</span></div><div class="line"><a name="l00029"></a><span class="lineno">   29</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00030"></a><span class="lineno">   30</span>&#160;</div><div class="line"><a name="l00031"></a><span class="lineno">   31</span>&#160;<span class="preprocessor">#ifndef DATA_TYPE_MAX</span></div><div class="line"><a name="l00032"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#afa6b114aa46a23bb226c5c9547f9f623">   32</a></span>&#160;<span class="preprocessor">#define DATA_TYPE_MAX 0xFF</span></div><div class="line"><a name="l00033"></a><span class="lineno">   33</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00034"></a><span class="lineno">   34</span>&#160;</div><div class="line"><a name="l00035"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">   35</a></span>&#160;__constant <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16) <a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a> = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16))(<a class="code" href="minmaxloc_8cl.xhtml#a17decce255d6f94a60379a61a54e3bac">DATA_TYPE_MIN</a>);</div><div class="line"><a name="l00036"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">   36</a></span>&#160;__constant <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16) <a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a> = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16))(<a class="code" href="minmaxloc_8cl.xhtml#afa6b114aa46a23bb226c5c9547f9f623">DATA_TYPE_MAX</a>);</div><div class="line"><a name="l00037"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">   37</a></span>&#160;__constant uint16 <a class="code" href="minmaxloc_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a> = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);</div><div class="line"><a name="l00038"></a><span class="lineno">   38</span>&#160;</div><div class="line"><a name="l00054"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a57351025b140b0b9ef881714d32d84d3">   54</a></span>&#160;__kernel <span class="keywordtype">void</span> <a class="code" href="minmaxloc_8cl.xhtml#a57351025b140b0b9ef881714d32d84d3">minmax</a>(</div><div class="line"><a name="l00055"></a><span class="lineno">   55</span>&#160;    <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src),</div><div class="line"><a name="l00056"></a><span class="lineno">   56</span>&#160;    __global <span class="keywordtype">int</span> *min_max,</div><div class="line"><a name="l00057"></a><span class="lineno">   57</span>&#160;    uint          width)</div><div class="line"><a name="l00058"></a><span class="lineno">   58</span>&#160;{</div><div class="line"><a name="l00059"></a><span class="lineno">   59</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00060"></a><span class="lineno">   60</span>&#160;</div><div class="line"><a name="l00061"></a><span class="lineno">   61</span>&#160;    <span class="comment">// Initialize local minimum and local maximum</span></div><div class="line"><a name="l00062"></a><span class="lineno">   62</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00063"></a><span class="lineno">   63</span>&#160;    local_min = <a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>;</div><div class="line"><a name="l00064"></a><span class="lineno">   64</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00065"></a><span class="lineno">   65</span>&#160;    local_max = <a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>;</div><div class="line"><a name="l00066"></a><span class="lineno">   66</span>&#160;</div><div class="line"><a name="l00067"></a><span class="lineno">   67</span>&#160;    <span class="comment">// Calculate min/max of row</span></div><div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;    uint width4 = width &gt;&gt; 4;</div><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;    <span class="keywordflow">for</span>(uint i = 0; i &lt; width4; i++)</div><div class="line"><a name="l00070"></a><span class="lineno">   70</span>&#160;    {</div><div class="line"><a name="l00071"></a><span class="lineno">   71</span>&#160;        <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00072"></a><span class="lineno">   72</span>&#160;        data      = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; 4, 0));</div><div class="line"><a name="l00073"></a><span class="lineno">   73</span>&#160;        local_min = min(data, local_min);</div><div class="line"><a name="l00074"></a><span class="lineno">   74</span>&#160;        local_max = max(data, local_max);</div><div class="line"><a name="l00075"></a><span class="lineno">   75</span>&#160;    }</div><div class="line"><a name="l00076"></a><span class="lineno">   76</span>&#160;</div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00078"></a><span class="lineno">   78</span>&#160;    <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00079"></a><span class="lineno">   79</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00080"></a><span class="lineno">   80</span>&#160;    data = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width4 &lt;&lt; 4, 0));</div><div class="line"><a name="l00081"></a><span class="lineno">   81</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00082"></a><span class="lineno">   82</span>&#160;    widx      = <a class="code" href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint16)(width4 &lt;&lt; 4) + <a class="code" href="minmaxloc_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a>) &lt; width, <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16));</div><div class="line"><a name="l00083"></a><span class="lineno">   83</span>&#160;    local_max = max(local_max, select(<a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>, data, widx));</div><div class="line"><a name="l00084"></a><span class="lineno">   84</span>&#160;    local_min = min(local_min, select(<a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>, data, widx));</div><div class="line"><a name="l00085"></a><span class="lineno">   85</span>&#160;<span class="preprocessor">#endif</span></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;    <span class="comment">// Perform min/max reduction</span></div><div class="line"><a name="l00088"></a><span class="lineno">   88</span>&#160;    local_min.s01234567 = min(local_min.s01234567, local_min.s89ABCDEF);</div><div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    local_max.s01234567 = max(local_max.s01234567, local_max.s89ABCDEF);</div><div class="line"><a name="l00090"></a><span class="lineno">   90</span>&#160;</div><div class="line"><a name="l00091"></a><span class="lineno">   91</span>&#160;    local_min.s0123 = min(local_min.s0123, local_min.s4567);</div><div class="line"><a name="l00092"></a><span class="lineno">   92</span>&#160;    local_max.s0123 = max(local_max.s0123, local_max.s4567);</div><div class="line"><a name="l00093"></a><span class="lineno">   93</span>&#160;</div><div class="line"><a name="l00094"></a><span class="lineno">   94</span>&#160;    local_min.s01 = min(local_min.s01, local_min.s23);</div><div class="line"><a name="l00095"></a><span class="lineno">   95</span>&#160;    local_max.s01 = max(local_max.s01, local_max.s23);</div><div class="line"><a name="l00096"></a><span class="lineno">   96</span>&#160;</div><div class="line"><a name="l00097"></a><span class="lineno">   97</span>&#160;    local_min.s0 = min(local_min.s0, local_min.s1);</div><div class="line"><a name="l00098"></a><span class="lineno">   98</span>&#160;    local_max.s0 = max(local_max.s0, local_max.s1);</div><div class="line"><a name="l00099"></a><span class="lineno">   99</span>&#160;</div><div class="line"><a name="l00100"></a><span class="lineno">  100</span>&#160;    <span class="comment">// Update global min/max</span></div><div class="line"><a name="l00101"></a><span class="lineno">  101</span>&#160;    atomic_min(&amp;min_max[0], local_min.s0);</div><div class="line"><a name="l00102"></a><span class="lineno">  102</span>&#160;    atomic_max(&amp;min_max[1], local_max.s0);</div><div class="line"><a name="l00103"></a><span class="lineno">  103</span>&#160;}</div><div class="line"><a name="l00104"></a><span class="lineno">  104</span>&#160;</div><div class="line"><a name="l00123"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a31ce1a75df24b2037dc76f1748236579">  123</a></span>&#160;__kernel <span class="keywordtype">void</span> <a class="code" href="minmaxloc_8cl.xhtml#a31ce1a75df24b2037dc76f1748236579">minmaxloc</a>(</div><div class="line"><a name="l00124"></a><span class="lineno">  124</span>&#160;    <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src),</div><div class="line"><a name="l00125"></a><span class="lineno">  125</span>&#160;    __global <span class="keywordtype">int</span> *min_max,</div><div class="line"><a name="l00126"></a><span class="lineno">  126</span>&#160;    __global uint *min_max_count</div><div class="line"><a name="l00127"></a><span class="lineno">  127</span>&#160;#<span class="keywordflow">if</span> defined        LOCATE_MIN</div><div class="line"><a name="l00128"></a><span class="lineno">  128</span>&#160;    ,</div><div class="line"><a name="l00129"></a><span class="lineno">  129</span>&#160;    __global <a class="code" href="struct_coordinates2_d.xhtml">Coordinates2D</a> *min_loc, uint max_min_loc_count</div><div class="line"><a name="l00130"></a><span class="lineno">  130</span>&#160;#endif</div><div class="line"><a name="l00131"></a><span class="lineno">  131</span>&#160;#<span class="keywordflow">if</span> defined LOCATE_MAX</div><div class="line"><a name="l00132"></a><span class="lineno">  132</span>&#160;    ,</div><div class="line"><a name="l00133"></a><span class="lineno">  133</span>&#160;    __global <a class="code" href="struct_coordinates2_d.xhtml">Coordinates2D</a> *max_loc, uint max_max_loc_count</div><div class="line"><a name="l00134"></a><span class="lineno">  134</span>&#160;#endif</div><div class="line"><a name="l00135"></a><span class="lineno">  135</span>&#160;)</div><div class="line"><a name="l00136"></a><span class="lineno">  136</span>&#160;{</div><div class="line"><a name="l00137"></a><span class="lineno">  137</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00138"></a><span class="lineno">  138</span>&#160;</div><div class="line"><a name="l00139"></a><span class="lineno">  139</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> *)src.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00140"></a><span class="lineno">  140</span>&#160;<span class="preprocessor">#if defined COUNT_MIN_MAX</span></div><div class="line"><a name="l00141"></a><span class="lineno">  141</span>&#160;    <span class="keywordflow">if</span>(value == min_max[0])</div><div class="line"><a name="l00142"></a><span class="lineno">  142</span>&#160;    {</div><div class="line"><a name="l00143"></a><span class="lineno">  143</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[0]);</div><div class="line"><a name="l00144"></a><span class="lineno">  144</span>&#160;<span class="preprocessor">#if defined  LOCATE_MIN</span></div><div class="line"><a name="l00145"></a><span class="lineno">  145</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_min_loc_count)</div><div class="line"><a name="l00146"></a><span class="lineno">  146</span>&#160;        {</div><div class="line"><a name="l00147"></a><span class="lineno">  147</span>&#160;            min_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00148"></a><span class="lineno">  148</span>&#160;            min_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00149"></a><span class="lineno">  149</span>&#160;        }</div><div class="line"><a name="l00150"></a><span class="lineno">  150</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00151"></a><span class="lineno">  151</span>&#160;    }</div><div class="line"><a name="l00152"></a><span class="lineno">  152</span>&#160;    <span class="keywordflow">if</span>(value == min_max[1])</div><div class="line"><a name="l00153"></a><span class="lineno">  153</span>&#160;    {</div><div class="line"><a name="l00154"></a><span class="lineno">  154</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[1]);</div><div class="line"><a name="l00155"></a><span class="lineno">  155</span>&#160;<span class="preprocessor">#if defined  LOCATE_MAX</span></div><div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_max_loc_count)</div><div class="line"><a name="l00157"></a><span class="lineno">  157</span>&#160;        {</div><div class="line"><a name="l00158"></a><span class="lineno">  158</span>&#160;            max_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00159"></a><span class="lineno">  159</span>&#160;            max_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00160"></a><span class="lineno">  160</span>&#160;        }</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00162"></a><span class="lineno">  162</span>&#160;    }</div><div class="line"><a name="l00163"></a><span class="lineno">  163</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00164"></a><span class="lineno">  164</span>&#160;}</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>
+<a href="minmaxloc_8cl.xhtml">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno">    1</span>&#160;<span class="comment">/*</span></div><div class="line"><a name="l00002"></a><span class="lineno">    2</span>&#160;<span class="comment"> * Copyright (c) 2017 ARM Limited.</span></div><div class="line"><a name="l00003"></a><span class="lineno">    3</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno">    4</span>&#160;<span class="comment"> * SPDX-License-Identifier: MIT</span></div><div class="line"><a name="l00005"></a><span class="lineno">    5</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00006"></a><span class="lineno">    6</span>&#160;<span class="comment"> * Permission is hereby granted, free of charge, to any person obtaining a copy</span></div><div class="line"><a name="l00007"></a><span class="lineno">    7</span>&#160;<span class="comment"> * of this software and associated documentation files (the &quot;Software&quot;), to</span></div><div class="line"><a name="l00008"></a><span class="lineno">    8</span>&#160;<span class="comment"> * deal in the Software without restriction, including without limitation the</span></div><div class="line"><a name="l00009"></a><span class="lineno">    9</span>&#160;<span class="comment"> * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or</span></div><div class="line"><a name="l00010"></a><span class="lineno">   10</span>&#160;<span class="comment"> * sell copies of the Software, and to permit persons to whom the Software is</span></div><div class="line"><a name="l00011"></a><span class="lineno">   11</span>&#160;<span class="comment"> * furnished to do so, subject to the following conditions:</span></div><div class="line"><a name="l00012"></a><span class="lineno">   12</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00013"></a><span class="lineno">   13</span>&#160;<span class="comment"> * The above copyright notice and this permission notice shall be included in all</span></div><div class="line"><a name="l00014"></a><span class="lineno">   14</span>&#160;<span class="comment"> * copies or substantial portions of the Software.</span></div><div class="line"><a name="l00015"></a><span class="lineno">   15</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00016"></a><span class="lineno">   16</span>&#160;<span class="comment"> * THE SOFTWARE IS PROVIDED &quot;AS IS&quot;, WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</span></div><div class="line"><a name="l00017"></a><span class="lineno">   17</span>&#160;<span class="comment"> * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,</span></div><div class="line"><a name="l00018"></a><span class="lineno">   18</span>&#160;<span class="comment"> * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE</span></div><div class="line"><a name="l00019"></a><span class="lineno">   19</span>&#160;<span class="comment"> * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER</span></div><div class="line"><a name="l00020"></a><span class="lineno">   20</span>&#160;<span class="comment"> * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,</span></div><div class="line"><a name="l00021"></a><span class="lineno">   21</span>&#160;<span class="comment"> * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE</span></div><div class="line"><a name="l00022"></a><span class="lineno">   22</span>&#160;<span class="comment"> * SOFTWARE.</span></div><div class="line"><a name="l00023"></a><span class="lineno">   23</span>&#160;<span class="comment"> */</span></div><div class="line"><a name="l00024"></a><span class="lineno">   24</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="helpers_8h.xhtml">helpers.h</a>&quot;</span></div><div class="line"><a name="l00025"></a><span class="lineno">   25</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="types_8h.xhtml">types.h</a>&quot;</span></div><div class="line"><a name="l00026"></a><span class="lineno">   26</span>&#160;</div><div class="line"><a name="l00027"></a><span class="lineno">   27</span>&#160;<span class="preprocessor">#ifndef DATA_TYPE_MIN</span></div><div class="line"><a name="l00028"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a17decce255d6f94a60379a61a54e3bac">   28</a></span>&#160;<span class="preprocessor">#define DATA_TYPE_MIN 0x0</span></div><div class="line"><a name="l00029"></a><span class="lineno">   29</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00030"></a><span class="lineno">   30</span>&#160;</div><div class="line"><a name="l00031"></a><span class="lineno">   31</span>&#160;<span class="preprocessor">#ifndef DATA_TYPE_MAX</span></div><div class="line"><a name="l00032"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#afa6b114aa46a23bb226c5c9547f9f623">   32</a></span>&#160;<span class="preprocessor">#define DATA_TYPE_MAX 0xFF</span></div><div class="line"><a name="l00033"></a><span class="lineno">   33</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00034"></a><span class="lineno">   34</span>&#160;</div><div class="line"><a name="l00035"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">   35</a></span>&#160;__constant <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16) <a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a> = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16))(<a class="code" href="minmaxloc_8cl.xhtml#a17decce255d6f94a60379a61a54e3bac">DATA_TYPE_MIN</a>);</div><div class="line"><a name="l00036"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">   36</a></span>&#160;__constant <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16) <a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a> = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16))(<a class="code" href="minmaxloc_8cl.xhtml#afa6b114aa46a23bb226c5c9547f9f623">DATA_TYPE_MAX</a>);</div><div class="line"><a name="l00037"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">   37</a></span>&#160;__constant uint16 <a class="code" href="minmaxloc_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a> = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);</div><div class="line"><a name="l00038"></a><span class="lineno">   38</span>&#160;</div><div class="line"><a name="l00054"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a57351025b140b0b9ef881714d32d84d3">   54</a></span>&#160;__kernel <span class="keywordtype">void</span> <a class="code" href="minmaxloc_8cl.xhtml#a57351025b140b0b9ef881714d32d84d3">minmax</a>(</div><div class="line"><a name="l00055"></a><span class="lineno">   55</span>&#160;    <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src),</div><div class="line"><a name="l00056"></a><span class="lineno">   56</span>&#160;    __global <span class="keywordtype">int</span> *min_max,</div><div class="line"><a name="l00057"></a><span class="lineno">   57</span>&#160;    uint          width)</div><div class="line"><a name="l00058"></a><span class="lineno">   58</span>&#160;{</div><div class="line"><a name="l00059"></a><span class="lineno">   59</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00060"></a><span class="lineno">   60</span>&#160;</div><div class="line"><a name="l00061"></a><span class="lineno">   61</span>&#160;    <span class="comment">// Initialize local minimum and local maximum</span></div><div class="line"><a name="l00062"></a><span class="lineno">   62</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00063"></a><span class="lineno">   63</span>&#160;    local_min = <a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>;</div><div class="line"><a name="l00064"></a><span class="lineno">   64</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00065"></a><span class="lineno">   65</span>&#160;    local_max = <a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>;</div><div class="line"><a name="l00066"></a><span class="lineno">   66</span>&#160;</div><div class="line"><a name="l00067"></a><span class="lineno">   67</span>&#160;    <span class="comment">// Calculate min/max of row</span></div><div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;    uint width4 = width &gt;&gt; 4;</div><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;    <span class="keywordflow">for</span>(uint i = 0; i &lt; width4; i++)</div><div class="line"><a name="l00070"></a><span class="lineno">   70</span>&#160;    {</div><div class="line"><a name="l00071"></a><span class="lineno">   71</span>&#160;        <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00072"></a><span class="lineno">   72</span>&#160;        data      = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; 4, 0));</div><div class="line"><a name="l00073"></a><span class="lineno">   73</span>&#160;        local_min = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#aabcf39e3917f842dbc5fbb0d802f24d5">min</a>(data, local_min);</div><div class="line"><a name="l00074"></a><span class="lineno">   74</span>&#160;        local_max = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">max</a>(data, local_max);</div><div class="line"><a name="l00075"></a><span class="lineno">   75</span>&#160;    }</div><div class="line"><a name="l00076"></a><span class="lineno">   76</span>&#160;</div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00078"></a><span class="lineno">   78</span>&#160;    <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00079"></a><span class="lineno">   79</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00080"></a><span class="lineno">   80</span>&#160;    data = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width4 &lt;&lt; 4, 0));</div><div class="line"><a name="l00081"></a><span class="lineno">   81</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00082"></a><span class="lineno">   82</span>&#160;    widx      = <a class="code" href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint16)(width4 &lt;&lt; 4) + <a class="code" href="minmaxloc_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a>) &lt; width, <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16));</div><div class="line"><a name="l00083"></a><span class="lineno">   83</span>&#160;    local_max = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">max</a>(local_max, select(<a class="code" href="minmaxloc_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>, data, widx));</div><div class="line"><a name="l00084"></a><span class="lineno">   84</span>&#160;    local_min = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#aabcf39e3917f842dbc5fbb0d802f24d5">min</a>(local_min, select(<a class="code" href="minmaxloc_8cl.xhtml#a4464d6f922ea17b4a9ca6a2cec7ddb75">type_max</a>, data, widx));</div><div class="line"><a name="l00085"></a><span class="lineno">   85</span>&#160;<span class="preprocessor">#endif</span></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;    <span class="comment">// Perform min/max reduction</span></div><div class="line"><a name="l00088"></a><span class="lineno">   88</span>&#160;    local_min.s01234567 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#aabcf39e3917f842dbc5fbb0d802f24d5">min</a>(local_min.s01234567, local_min.s89ABCDEF);</div><div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    local_max.s01234567 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">max</a>(local_max.s01234567, local_max.s89ABCDEF);</div><div class="line"><a name="l00090"></a><span class="lineno">   90</span>&#160;</div><div class="line"><a name="l00091"></a><span class="lineno">   91</span>&#160;    local_min.s0123 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#aabcf39e3917f842dbc5fbb0d802f24d5">min</a>(local_min.s0123, local_min.s4567);</div><div class="line"><a name="l00092"></a><span class="lineno">   92</span>&#160;    local_max.s0123 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">max</a>(local_max.s0123, local_max.s4567);</div><div class="line"><a name="l00093"></a><span class="lineno">   93</span>&#160;</div><div class="line"><a name="l00094"></a><span class="lineno">   94</span>&#160;    local_min.s01 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#aabcf39e3917f842dbc5fbb0d802f24d5">min</a>(local_min.s01, local_min.s23);</div><div class="line"><a name="l00095"></a><span class="lineno">   95</span>&#160;    local_max.s01 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">max</a>(local_max.s01, local_max.s23);</div><div class="line"><a name="l00096"></a><span class="lineno">   96</span>&#160;</div><div class="line"><a name="l00097"></a><span class="lineno">   97</span>&#160;    local_min.s0 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#aabcf39e3917f842dbc5fbb0d802f24d5">min</a>(local_min.s0, local_min.s1);</div><div class="line"><a name="l00098"></a><span class="lineno">   98</span>&#160;    local_max.s0 = <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">max</a>(local_max.s0, local_max.s1);</div><div class="line"><a name="l00099"></a><span class="lineno">   99</span>&#160;</div><div class="line"><a name="l00100"></a><span class="lineno">  100</span>&#160;    <span class="comment">// Update global min/max</span></div><div class="line"><a name="l00101"></a><span class="lineno">  101</span>&#160;    atomic_min(&amp;min_max[0], local_min.s0);</div><div class="line"><a name="l00102"></a><span class="lineno">  102</span>&#160;    atomic_max(&amp;min_max[1], local_max.s0);</div><div class="line"><a name="l00103"></a><span class="lineno">  103</span>&#160;}</div><div class="line"><a name="l00104"></a><span class="lineno">  104</span>&#160;</div><div class="line"><a name="l00123"></a><span class="lineno"><a class="line" href="minmaxloc_8cl.xhtml#a31ce1a75df24b2037dc76f1748236579">  123</a></span>&#160;__kernel <span class="keywordtype">void</span> <a class="code" href="minmaxloc_8cl.xhtml#a31ce1a75df24b2037dc76f1748236579">minmaxloc</a>(</div><div class="line"><a name="l00124"></a><span class="lineno">  124</span>&#160;    <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src),</div><div class="line"><a name="l00125"></a><span class="lineno">  125</span>&#160;    __global <span class="keywordtype">int</span> *min_max,</div><div class="line"><a name="l00126"></a><span class="lineno">  126</span>&#160;    __global uint *min_max_count</div><div class="line"><a name="l00127"></a><span class="lineno">  127</span>&#160;#<span class="keywordflow">if</span> defined        LOCATE_MIN</div><div class="line"><a name="l00128"></a><span class="lineno">  128</span>&#160;    ,</div><div class="line"><a name="l00129"></a><span class="lineno">  129</span>&#160;    __global <a class="code" href="struct_coordinates2_d.xhtml">Coordinates2D</a> *min_loc, uint max_min_loc_count</div><div class="line"><a name="l00130"></a><span class="lineno">  130</span>&#160;#endif</div><div class="line"><a name="l00131"></a><span class="lineno">  131</span>&#160;#<span class="keywordflow">if</span> defined LOCATE_MAX</div><div class="line"><a name="l00132"></a><span class="lineno">  132</span>&#160;    ,</div><div class="line"><a name="l00133"></a><span class="lineno">  133</span>&#160;    __global <a class="code" href="struct_coordinates2_d.xhtml">Coordinates2D</a> *max_loc, uint max_max_loc_count</div><div class="line"><a name="l00134"></a><span class="lineno">  134</span>&#160;#endif</div><div class="line"><a name="l00135"></a><span class="lineno">  135</span>&#160;)</div><div class="line"><a name="l00136"></a><span class="lineno">  136</span>&#160;{</div><div class="line"><a name="l00137"></a><span class="lineno">  137</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00138"></a><span class="lineno">  138</span>&#160;</div><div class="line"><a name="l00139"></a><span class="lineno">  139</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> *)src.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00140"></a><span class="lineno">  140</span>&#160;<span class="preprocessor">#if defined COUNT_MIN_MAX</span></div><div class="line"><a name="l00141"></a><span class="lineno">  141</span>&#160;    <span class="keywordflow">if</span>(value == min_max[0])</div><div class="line"><a name="l00142"></a><span class="lineno">  142</span>&#160;    {</div><div class="line"><a name="l00143"></a><span class="lineno">  143</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[0]);</div><div class="line"><a name="l00144"></a><span class="lineno">  144</span>&#160;<span class="preprocessor">#if defined  LOCATE_MIN</span></div><div class="line"><a name="l00145"></a><span class="lineno">  145</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_min_loc_count)</div><div class="line"><a name="l00146"></a><span class="lineno">  146</span>&#160;        {</div><div class="line"><a name="l00147"></a><span class="lineno">  147</span>&#160;            min_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00148"></a><span class="lineno">  148</span>&#160;            min_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00149"></a><span class="lineno">  149</span>&#160;        }</div><div class="line"><a name="l00150"></a><span class="lineno">  150</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00151"></a><span class="lineno">  151</span>&#160;    }</div><div class="line"><a name="l00152"></a><span class="lineno">  152</span>&#160;    <span class="keywordflow">if</span>(value == min_max[1])</div><div class="line"><a name="l00153"></a><span class="lineno">  153</span>&#160;    {</div><div class="line"><a name="l00154"></a><span class="lineno">  154</span>&#160;        uint idx = atomic_inc(&amp;min_max_count[1]);</div><div class="line"><a name="l00155"></a><span class="lineno">  155</span>&#160;<span class="preprocessor">#if defined  LOCATE_MAX</span></div><div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;        <span class="keywordflow">if</span>(idx &lt; max_max_loc_count)</div><div class="line"><a name="l00157"></a><span class="lineno">  157</span>&#160;        {</div><div class="line"><a name="l00158"></a><span class="lineno">  158</span>&#160;            max_loc[idx].x = get_global_id(0);</div><div class="line"><a name="l00159"></a><span class="lineno">  159</span>&#160;            max_loc[idx].y = get_global_id(1);</div><div class="line"><a name="l00160"></a><span class="lineno">  160</span>&#160;        }</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00162"></a><span class="lineno">  162</span>&#160;    }</div><div class="line"><a name="l00163"></a><span class="lineno">  163</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00164"></a><span class="lineno">  164</span>&#160;}</div><div class="ttc" id="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail_xhtml_aabcf39e3917f842dbc5fbb0d802f24d5"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#aabcf39e3917f842dbc5fbb0d802f24d5">arm_compute::test::fixed_point_arithmetic::detail::min</a></div><div class="ttdeci">fixed_point&lt; T &gt; min(fixed_point&lt; T &gt; x, fixed_point&lt; T &gt; y)</div><div class="ttdef"><b>Definition:</b> <a href="tests_2validation_2_fixed_point_8h_source.xhtml#l00875">FixedPoint.h:875</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="minmaxloc_8cl_xhtml_afa6b114aa46a23bb226c5c9547f9f623"><div class="ttname"><a href="minmaxloc_8cl.xhtml#afa6b114aa46a23bb226c5c9547f9f623">DATA_TYPE_MAX</a></div><div class="ttdeci">#define DATA_TYPE_MAX</div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00032">minmaxloc.cl:32</a></div></div>
 <div class="ttc" id="struct_coordinates2_d_xhtml"><div class="ttname"><a href="struct_coordinates2_d.xhtml">Coordinates2D</a></div><div class="ttdoc">2D Coordinates structure </div><div class="ttdef"><b>Definition:</b> <a href="types_8h_source.xhtml#l00028">types.h:28</a></div></div>
 <div class="ttc" id="helpers_8h_xhtml_a22f42fcf2077d951271df83b55c1a71a"><div class="ttname"><a href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a></div><div class="ttdeci">#define IMAGE_DECLARATION(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00049">helpers.h:49</a></div></div>
@@ -131,6 +133,7 @@
 <div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00097">helpers.h:97</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#l00036">minmaxloc.cl:36</a></div></div>
 <div class="ttc" id="minmaxloc_8cl_xhtml_a0712735973f172ac9efc7d48a31e47ad"><div class="ttname"><a href="minmaxloc_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a></div><div class="ttdeci">__constant uint16 idx16</div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00037">minmaxloc.cl:37</a></div></div>
+<div class="ttc" id="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail_xhtml_ad91bb73431b4de1f4946ed949d444849"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">arm_compute::test::fixed_point_arithmetic::detail::max</a></div><div class="ttdeci">fixed_point&lt; T &gt; max(fixed_point&lt; T &gt; x, fixed_point&lt; T &gt; y)</div><div class="ttdef"><b>Definition:</b> <a href="tests_2validation_2_fixed_point_8h_source.xhtml#l00880">FixedPoint.h:880</a></div></div>
 <div class="ttc" id="helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00035">helpers.h:35</a></div></div>
 <div class="ttc" id="minmaxloc_8cl_xhtml_a17decce255d6f94a60379a61a54e3bac"><div class="ttname"><a href="minmaxloc_8cl.xhtml#a17decce255d6f94a60379a61a54e3bac">DATA_TYPE_MIN</a></div><div class="ttdeci">#define DATA_TYPE_MIN</div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00028">minmaxloc.cl:28</a></div></div>
 <div class="ttc" id="minmaxloc_8cl_xhtml_a57351025b140b0b9ef881714d32d84d3"><div class="ttname"><a href="minmaxloc_8cl.xhtml#a57351025b140b0b9ef881714d32d84d3">minmax</a></div><div class="ttdeci">__kernel void minmax(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global int *min_max, uint width)</div><div class="ttdoc">This function identifies the min and maximum value of an input image. </div><div class="ttdef"><b>Definition:</b> <a href="minmaxloc_8cl_source.xhtml#l00054">minmaxloc.cl:54</a></div></div>
@@ -140,7 +143,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 Wed May 3 2017 17:20:04 for ARM Compute Library by
+    <li class="footer">Generated on Fri Jun 23 2017 15:44:33 for Compute Library by
     <a href="http://www.doxygen.org/index.html">
     <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
   </ul>