arm_compute v18.08
diff --git a/documentation/reduction__operation_8cl_source.xhtml b/documentation/reduction__operation_8cl_source.xhtml
index 7dbee54..7fd37ac 100644
--- a/documentation/reduction__operation_8cl_source.xhtml
+++ b/documentation/reduction__operation_8cl_source.xhtml
@@ -40,7 +40,7 @@
<tr style="height: 56px;">
<td style="padding-left: 0.5em;">
<div id="projectname">Compute Library
-  <span id="projectnumber">18.05</span>
+  <span id="projectnumber">18.08</span>
</div>
</td>
</tr>
@@ -117,16 +117,17 @@
<div class="title">reduction_operation.cl</div> </div>
</div><!--header-->
<div class="contents">
-<a href="reduction__operation_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> <span class="comment">/*</span></div><div class="line"><a name="l00002"></a><span class="lineno"> 2</span> <span class="comment"> * Copyright (c) 2016-2018 ARM Limited.</span></div><div class="line"><a name="l00003"></a><span class="lineno"> 3</span> <span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno"> 4</span> <span class="comment"> * SPDX-License-Identifier: MIT</span></div><div class="line"><a name="l00005"></a><span class="lineno"> 5</span> <span class="comment"> *</span></div><div class="line"><a name="l00006"></a><span class="lineno"> 6</span> <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> <span class="comment"> * of this software and associated documentation files (the "Software"), to</span></div><div class="line"><a name="l00008"></a><span class="lineno"> 8</span> <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> <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> <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> <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> <span class="comment"> *</span></div><div class="line"><a name="l00013"></a><span class="lineno"> 13</span> <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> <span class="comment"> * copies or substantial portions of the Software.</span></div><div class="line"><a name="l00015"></a><span class="lineno"> 15</span> <span class="comment"> *</span></div><div class="line"><a name="l00016"></a><span class="lineno"> 16</span> <span class="comment"> * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</span></div><div class="line"><a name="l00017"></a><span class="lineno"> 17</span> <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> <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> <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> <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> <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> <span class="comment"> * SOFTWARE.</span></div><div class="line"><a name="l00023"></a><span class="lineno"> 23</span> <span class="comment"> */</span></div><div class="line"><a name="l00024"></a><span class="lineno"> 24</span> <span class="preprocessor">#include "<a class="code" href="helpers_8h.xhtml">helpers.h</a>"</span></div><div class="line"><a name="l00025"></a><span class="lineno"> 25</span> </div><div class="line"><a name="l00032"></a><span class="lineno"><a class="line" href="reduction__operation_8cl.xhtml#a48cc0583beba45dfba6d94fb6dc3c59a"> 32</a></span> <span class="keyword">inline</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> <a class="code" href="reduction__operation_8cl.xhtml#a48cc0583beba45dfba6d94fb6dc3c59a">square_sum</a>(__global <span class="keyword">const</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *input)</div><div class="line"><a name="l00033"></a><span class="lineno"> 33</span> {</div><div class="line"><a name="l00034"></a><span class="lineno"> 34</span>  <a class="code" href="fixed__point_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="l00035"></a><span class="lineno"> 35</span>  in = vload16(0, input);</div><div class="line"><a name="l00036"></a><span class="lineno"> 36</span> </div><div class="line"><a name="l00037"></a><span class="lineno"> 37</span>  in *= in;</div><div class="line"><a name="l00038"></a><span class="lineno"> 38</span> </div><div class="line"><a name="l00039"></a><span class="lineno"> 39</span>  in.s01234567 += in.s89ABCDEF;</div><div class="line"><a name="l00040"></a><span class="lineno"> 40</span>  in.s0123 += in.s4567;</div><div class="line"><a name="l00041"></a><span class="lineno"> 41</span>  in.s01 += in.s23;</div><div class="line"><a name="l00042"></a><span class="lineno"> 42</span> </div><div class="line"><a name="l00043"></a><span class="lineno"> 43</span>  <span class="keywordflow">return</span> (in.s0 + in.s1);</div><div class="line"><a name="l00044"></a><span class="lineno"> 44</span> }</div><div class="line"><a name="l00045"></a><span class="lineno"> 45</span> </div><div class="line"><a name="l00052"></a><span class="lineno"><a class="line" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782"> 52</a></span> <span class="keyword">inline</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> <a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>(__global <span class="keyword">const</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *input)</div><div class="line"><a name="l00053"></a><span class="lineno"> 53</span> {</div><div class="line"><a name="l00054"></a><span class="lineno"> 54</span>  <a class="code" href="fixed__point_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="l00055"></a><span class="lineno"> 55</span>  in = vload16(0, input);</div><div class="line"><a name="l00056"></a><span class="lineno"> 56</span> </div><div class="line"><a name="l00057"></a><span class="lineno"> 57</span>  in.s01234567 += in.s89ABCDEF;</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>  in.s0123 += in.s4567;</div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>  in.s01 += in.s23;</div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span> </div><div class="line"><a name="l00061"></a><span class="lineno"> 61</span>  <span class="keywordflow">return</span> (in.s0 + in.s1);</div><div class="line"><a name="l00062"></a><span class="lineno"> 62</span> }</div><div class="line"><a name="l00063"></a><span class="lineno"> 63</span> </div><div class="line"><a name="l00084"></a><span class="lineno"><a class="line" href="reduction__operation_8cl.xhtml#a40c2cb11f583c331d7aefa2c61e6313f"> 84</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="reduction__operation_8cl.xhtml#a40c2cb11f583c331d7aefa2c61e6313f">reduction_operation</a>(</div><div class="line"><a name="l00085"></a><span class="lineno"> 85</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a>),</div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(partial_sum),</div><div class="line"><a name="l00087"></a><span class="lineno"> 87</span>  __local <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *local_sums)</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span> {</div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span>  <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>  <a class="code" href="struct_image.xhtml">Image</a> partial_sum = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(partial_sum);</div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span> </div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span>  <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> lsize = get_local_size(0);</div><div class="line"><a name="l00093"></a><span class="lineno"> 93</span>  <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> lid = get_local_id(0);</div><div class="line"><a name="l00094"></a><span class="lineno"> 94</span> </div><div class="line"><a name="l00095"></a><span class="lineno"> 95</span>  <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> y = 0; y < get_local_size(1); ++y)</div><div class="line"><a name="l00096"></a><span class="lineno"> 96</span>  {</div><div class="line"><a name="l00097"></a><span class="lineno"> 97</span>  local_sums[lid] = OPERATION((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, 0, y));</div><div class="line"><a name="l00098"></a><span class="lineno"> 98</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00099"></a><span class="lineno"> 99</span> </div><div class="line"><a name="l00100"></a><span class="lineno"> 100</span>  <span class="comment">// Perform parallel reduction</span></div><div class="line"><a name="l00101"></a><span class="lineno"> 101</span>  <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> i = lsize >> 1; i > 0; i >>= 1)</div><div class="line"><a name="l00102"></a><span class="lineno"> 102</span>  {</div><div class="line"><a name="l00103"></a><span class="lineno"> 103</span>  <span class="keywordflow">if</span>(lid < i)</div><div class="line"><a name="l00104"></a><span class="lineno"> 104</span>  {</div><div class="line"><a name="l00105"></a><span class="lineno"> 105</span>  local_sums[lid] += local_sums[lid + i];</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span>  }</div><div class="line"><a name="l00107"></a><span class="lineno"> 107</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00108"></a><span class="lineno"> 108</span>  }</div><div class="line"><a name="l00109"></a><span class="lineno"> 109</span> </div><div class="line"><a name="l00110"></a><span class="lineno"> 110</span>  <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00111"></a><span class="lineno"> 111</span>  {</div><div class="line"><a name="l00112"></a><span class="lineno"> 112</span>  ((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&partial_sum, get_group_id(0), y))[0] = local_sums[0];</div><div class="line"><a name="l00113"></a><span class="lineno"> 113</span>  }</div><div class="line"><a name="l00114"></a><span class="lineno"> 114</span>  }</div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span> }</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="reduction__operation_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> <span class="comment">/*</span></div><div class="line"><a name="l00002"></a><span class="lineno"> 2</span> <span class="comment"> * Copyright (c) 2016-2018 ARM Limited.</span></div><div class="line"><a name="l00003"></a><span class="lineno"> 3</span> <span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno"> 4</span> <span class="comment"> * SPDX-License-Identifier: MIT</span></div><div class="line"><a name="l00005"></a><span class="lineno"> 5</span> <span class="comment"> *</span></div><div class="line"><a name="l00006"></a><span class="lineno"> 6</span> <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> <span class="comment"> * of this software and associated documentation files (the "Software"), to</span></div><div class="line"><a name="l00008"></a><span class="lineno"> 8</span> <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> <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> <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> <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> <span class="comment"> *</span></div><div class="line"><a name="l00013"></a><span class="lineno"> 13</span> <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> <span class="comment"> * copies or substantial portions of the Software.</span></div><div class="line"><a name="l00015"></a><span class="lineno"> 15</span> <span class="comment"> *</span></div><div class="line"><a name="l00016"></a><span class="lineno"> 16</span> <span class="comment"> * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</span></div><div class="line"><a name="l00017"></a><span class="lineno"> 17</span> <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> <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> <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> <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> <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> <span class="comment"> * SOFTWARE.</span></div><div class="line"><a name="l00023"></a><span class="lineno"> 23</span> <span class="comment"> */</span></div><div class="line"><a name="l00024"></a><span class="lineno"> 24</span> <span class="preprocessor">#include "<a class="code" href="helpers_8h.xhtml">helpers.h</a>"</span></div><div class="line"><a name="l00025"></a><span class="lineno"> 25</span> </div><div class="line"><a name="l00032"></a><span class="lineno"><a class="line" href="reduction__operation_8cl.xhtml#a48cc0583beba45dfba6d94fb6dc3c59a"> 32</a></span> <span class="keyword">inline</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> <a class="code" href="reduction__operation_8cl.xhtml#a48cc0583beba45dfba6d94fb6dc3c59a">square_sum</a>(__global <span class="keyword">const</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *input)</div><div class="line"><a name="l00033"></a><span class="lineno"> 33</span> {</div><div class="line"><a name="l00034"></a><span class="lineno"> 34</span>  <a class="code" href="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="l00035"></a><span class="lineno"> 35</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a> = vload16(0, input);</div><div class="line"><a name="l00036"></a><span class="lineno"> 36</span> </div><div class="line"><a name="l00037"></a><span class="lineno"> 37</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a> *= <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>;</div><div class="line"><a name="l00038"></a><span class="lineno"> 38</span> </div><div class="line"><a name="l00039"></a><span class="lineno"> 39</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s01234567 += <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s89ABCDEF;</div><div class="line"><a name="l00040"></a><span class="lineno"> 40</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s0123 += <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s4567;</div><div class="line"><a name="l00041"></a><span class="lineno"> 41</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s01 += <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s23;</div><div class="line"><a name="l00042"></a><span class="lineno"> 42</span> </div><div class="line"><a name="l00043"></a><span class="lineno"> 43</span>  <span class="keywordflow">return</span> (<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s0 + <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s1);</div><div class="line"><a name="l00044"></a><span class="lineno"> 44</span> }</div><div class="line"><a name="l00045"></a><span class="lineno"> 45</span> </div><div class="line"><a name="l00052"></a><span class="lineno"><a class="line" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782"> 52</a></span> <span class="keyword">inline</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> <a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>(__global <span class="keyword">const</span> <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *input)</div><div class="line"><a name="l00053"></a><span class="lineno"> 53</span> {</div><div class="line"><a name="l00054"></a><span class="lineno"> 54</span>  <a class="code" href="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="l00055"></a><span class="lineno"> 55</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a> = vload16(0, input);</div><div class="line"><a name="l00056"></a><span class="lineno"> 56</span> </div><div class="line"><a name="l00057"></a><span class="lineno"> 57</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s01234567 += <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s89ABCDEF;</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s0123 += <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s4567;</div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s01 += <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s23;</div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span> </div><div class="line"><a name="l00061"></a><span class="lineno"> 61</span>  <span class="keywordflow">return</span> (<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s0 + <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a>.s1);</div><div class="line"><a name="l00062"></a><span class="lineno"> 62</span> }</div><div class="line"><a name="l00063"></a><span class="lineno"> 63</span> </div><div class="line"><a name="l00084"></a><span class="lineno"><a class="line" href="reduction__operation_8cl.xhtml#a40c2cb11f583c331d7aefa2c61e6313f"> 84</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="reduction__operation_8cl.xhtml#a40c2cb11f583c331d7aefa2c61e6313f">reduction_operation</a>(</div><div class="line"><a name="l00085"></a><span class="lineno"> 85</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a>),</div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(partial_sum),</div><div class="line"><a name="l00087"></a><span class="lineno"> 87</span>  __local <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *local_sums)</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span> {</div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span>  <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>  <a class="code" href="struct_image.xhtml">Image</a> partial_sum = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(partial_sum);</div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span> </div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span>  <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> lsize = get_local_size(0);</div><div class="line"><a name="l00093"></a><span class="lineno"> 93</span>  <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> lid = get_local_id(0);</div><div class="line"><a name="l00094"></a><span class="lineno"> 94</span> </div><div class="line"><a name="l00095"></a><span class="lineno"> 95</span>  <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> y = 0; y < get_local_size(1); ++y)</div><div class="line"><a name="l00096"></a><span class="lineno"> 96</span>  {</div><div class="line"><a name="l00097"></a><span class="lineno"> 97</span>  local_sums[lid] = OPERATION((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, 0, y));</div><div class="line"><a name="l00098"></a><span class="lineno"> 98</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00099"></a><span class="lineno"> 99</span> </div><div class="line"><a name="l00100"></a><span class="lineno"> 100</span>  <span class="comment">// Perform parallel reduction</span></div><div class="line"><a name="l00101"></a><span class="lineno"> 101</span>  <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> i = lsize >> 1; i > 0; i >>= 1)</div><div class="line"><a name="l00102"></a><span class="lineno"> 102</span>  {</div><div class="line"><a name="l00103"></a><span class="lineno"> 103</span>  <span class="keywordflow">if</span>(lid < i)</div><div class="line"><a name="l00104"></a><span class="lineno"> 104</span>  {</div><div class="line"><a name="l00105"></a><span class="lineno"> 105</span>  local_sums[lid] += local_sums[lid + i];</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span>  }</div><div class="line"><a name="l00107"></a><span class="lineno"> 107</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00108"></a><span class="lineno"> 108</span>  }</div><div class="line"><a name="l00109"></a><span class="lineno"> 109</span> </div><div class="line"><a name="l00110"></a><span class="lineno"> 110</span>  <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00111"></a><span class="lineno"> 111</span>  {</div><div class="line"><a name="l00112"></a><span class="lineno"> 112</span>  ((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&partial_sum, get_group_id(0), y))[0] = local_sums[0];</div><div class="line"><a name="l00113"></a><span class="lineno"> 113</span>  }</div><div class="line"><a name="l00114"></a><span class="lineno"> 114</span>  }</div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span> }</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="reduction__operation_8cl_xhtml_ab0df00f5333da51860deb93deb44a782"><div class="ttname"><a href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a></div><div class="ttdeci">DATA_TYPE sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate sum of a vector. </div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00052">reduction_operation.cl:52</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#l00068">helpers.h:68</a></div></div>
-<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00303">helpers.h:303</a></div></div>
+<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ad366eaa4cf2d106037d91c30795a5c76"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">arm_compute::test::validation::in</a></div><div class="ttdeci">CLTensor in</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_winograd_8cpp_source.xhtml#l00241">Winograd.cpp:241</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#l00076">helpers.h:76</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00311">helpers.h:311</a></div></div>
<div class="ttc" id="helpers_8h_xhtml"><div class="ttname"><a href="helpers_8h.xhtml">helpers.h</a></div></div>
-<div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00104">helpers.h:104</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="helpers_8h_source.xhtml#l00142">helpers.h:142</a></div></div>
-<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_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="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00112">helpers.h:112</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00056">helpers.h:56</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="helpers_8h_source.xhtml#l00150">helpers.h:150</a></div></div>
<div class="ttc" id="reduction__operation_8cl_xhtml_a40c2cb11f583c331d7aefa2c61e6313f"><div class="ttname"><a href="reduction__operation_8cl.xhtml#a40c2cb11f583c331d7aefa2c61e6313f">reduction_operation</a></div><div class="ttdeci">__kernel void reduction_operation(__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 uchar *partial_sum_ptr, uint partial_sum_stride_x, uint partial_sum_step_x, uint partial_sum_stride_y, uint partial_sum_step_y, uint partial_sum_offset_first_element_in_bytes, __local DATA_TYPE *local_sums)</div><div class="ttdoc">This kernel performs reduction given an operation. </div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00084">reduction_operation.cl:84</a></div></div>
-<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure & src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
+<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure & src</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
<div class="ttc" id="reduction__operation_8cl_xhtml_a48cc0583beba45dfba6d94fb6dc3c59a"><div class="ttname"><a href="reduction__operation_8cl.xhtml#a48cc0583beba45dfba6d94fb6dc3c59a">square_sum</a></div><div class="ttdeci">DATA_TYPE square_sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate square sum of a vector. </div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00032">reduction_operation.cl:32</a></div></div>
</div><!-- fragment --></div><!-- contents -->
</div><!-- doc-content -->
@@ -134,7 +135,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="reduction__operation_8cl.xhtml">reduction_operation.cl</a></li>
- <li class="footer">Generated on Wed May 23 2018 11:36:36 for Compute Library by
+ <li class="footer">Generated on Wed Aug 29 2018 15:31:29 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>