arm_compute v18.05
[platform/upstream/armcl.git] / documentation / _i_c_l_kernel_8h_source.xhtml
index bb8a5da..7aa354f 100644 (file)
@@ -4,7 +4,7 @@
 <head>
 <meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
 <meta http-equiv="X-UA-Compatible" content="IE=9"/>
-<meta name="generator" content="Doxygen 1.8.6"/>
+<meta name="generator" content="Doxygen 1.8.11"/>
 <meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
 <title>Compute Library: arm_compute/core/CL/ICLKernel.h Source File</title>
 <link href="tabs.css" rel="stylesheet" type="text/css"/>
 <script type="text/javascript" src="dynsections.js"></script>
 <link href="navtree.css" rel="stylesheet" type="text/css"/>
 <script type="text/javascript" src="resize.js"></script>
+<script type="text/javascript" src="navtreedata.js"></script>
 <script type="text/javascript" src="navtree.js"></script>
 <script type="text/javascript">
   $(document).ready(initResizable);
   $(window).load(resizeHeight);
 </script>
 <link href="search/search.css" rel="stylesheet" type="text/css"/>
+<script type="text/javascript" src="search/searchdata.js"></script>
 <script type="text/javascript" src="search/search.js"></script>
 <script type="text/javascript">
-  $(document).ready(function() { searchBox.OnSelectItem(0); });
+  $(document).ready(function() { init_search(); });
 </script>
 <script type="text/x-mathjax-config">
   MathJax.Hub.Config({
     extensions: ["tex2jax.js"],
     jax: ["input/TeX","output/HTML-CSS"],
 });
-</script><script src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
+</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
 <link href="doxygen.css" rel="stylesheet" type="text/css" />
 </head>
 <body>
@@ -38,7 +40,7 @@
  <tr style="height: 56px;">
   <td style="padding-left: 0.5em;">
    <div id="projectname">Compute Library
-   &#160;<span id="projectnumber">17.09</span>
+   &#160;<span id="projectnumber">18.05</span>
    </div>
   </td>
  </tr>
@@ -46,7 +48,7 @@
 </table>
 </div>
 <!-- end header part -->
-<!-- Generated by Doxygen 1.8.6 -->
+<!-- Generated by Doxygen 1.8.11 -->
 <script type="text/javascript">
 var searchBox = new SearchBox("searchBox", "search",false,'Search');
 </script>
@@ -101,7 +103,7 @@ $(document).ready(function(){initNavTree('_i_c_l_kernel_8h_source.xhtml','');});
      onmouseover="return searchBox.OnSearchSelectShow()"
      onmouseout="return searchBox.OnSearchSelectHide()"
      onkeydown="return searchBox.OnSearchSelectKey(event)">
-<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Data Structures</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Friends</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(10)"><span class="SelectionMark">&#160;</span>Macros</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(11)"><span class="SelectionMark">&#160;</span>Pages</a></div>
+</div>
 
 <!-- iframe showing the search results (closed by default) -->
 <div id="MSearchResultsWindow">
@@ -115,196 +117,58 @@ $(document).ready(function(){initNavTree('_i_c_l_kernel_8h_source.xhtml','');});
 <div class="title">ICLKernel.h</div>  </div>
 </div><!--header-->
 <div class="contents">
-<a href="_i_c_l_kernel_8h.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) 2016, 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">#ifndef __ARM_COMPUTE_ICLKERNEL_H__</span></div>
-<div class="line"><a name="l00025"></a><span class="lineno">   25</span>&#160;<span class="preprocessor"></span><span class="preprocessor">#define __ARM_COMPUTE_ICLKERNEL_H__</span></div>
-<div class="line"><a name="l00026"></a><span class="lineno">   26</span>&#160;<span class="preprocessor"></span></div>
-<div class="line"><a name="l00027"></a><span class="lineno">   27</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_c_l_kernel_library_8h.xhtml">arm_compute/core/CL/CLKernelLibrary.h</a>&quot;</span></div>
-<div class="line"><a name="l00028"></a><span class="lineno">   28</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_c_l_types_8h.xhtml">arm_compute/core/CL/CLTypes.h</a>&quot;</span></div>
-<div class="line"><a name="l00029"></a><span class="lineno">   29</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_open_c_l_8h.xhtml">arm_compute/core/CL/OpenCL.h</a>&quot;</span></div>
-<div class="line"><a name="l00030"></a><span class="lineno">   30</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_i_kernel_8h.xhtml">arm_compute/core/IKernel.h</a>&quot;</span></div>
-<div class="line"><a name="l00031"></a><span class="lineno">   31</span>&#160;</div>
-<div class="line"><a name="l00032"></a><span class="lineno">   32</span>&#160;<span class="preprocessor">#include &lt;string&gt;</span></div>
-<div class="line"><a name="l00033"></a><span class="lineno">   33</span>&#160;</div>
-<div class="line"><a name="l00034"></a><span class="lineno">   34</span>&#160;<span class="keyword">namespace </span>arm_compute</div>
-<div class="line"><a name="l00035"></a><span class="lineno">   35</span>&#160;{</div>
-<div class="line"><a name="l00036"></a><span class="lineno">   36</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
-<div class="line"><a name="l00037"></a><span class="lineno">   37</span>&#160;<span class="keyword">class </span>ICLArray;</div>
-<div class="line"><a name="l00038"></a><span class="lineno">   38</span>&#160;<span class="keyword">class </span><a class="code" href="namespacearm__compute.xhtml#af87af588094f04e537df98eaf0426cb6">ICLTensor</a>;</div>
-<div class="line"><a name="l00039"></a><span class="lineno">   39</span>&#160;<span class="keyword">class </span>Window;</div>
-<div class="line"><a name="l00040"></a><span class="lineno">   40</span>&#160;</div>
-<div class="line"><a name="l00042"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml">   42</a></span>&#160;<span class="keyword">class </span><a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml">ICLKernel</a> : <span class="keyword">public</span> <a class="code" href="classarm__compute_1_1_i_kernel.xhtml">IKernel</a></div>
-<div class="line"><a name="l00043"></a><span class="lineno">   43</span>&#160;{</div>
-<div class="line"><a name="l00044"></a><span class="lineno">   44</span>&#160;<span class="keyword">public</span>:</div>
-<div class="line"><a name="l00046"></a><span class="lineno">   46</span>&#160;    <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">ICLKernel</a>();</div>
-<div class="line"><a name="l00051"></a><span class="lineno">   51</span>&#160;    cl::Kernel &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">kernel</a>();</div>
-<div class="line"><a name="l00060"></a><span class="lineno">   60</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
-<div class="line"><a name="l00061"></a><span class="lineno">   61</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">add_1D_array_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div>
-<div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">add_1D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div>
-<div class="line"><a name="l00075"></a><span class="lineno">   75</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">add_2D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div>
-<div class="line"><a name="l00082"></a><span class="lineno">   82</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">add_3D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div>
-<div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">add_4D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div>
-<div class="line"><a name="l00094"></a><span class="lineno">   94</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a99fde125501fae43952af222c91236cd">num_arguments_per_1D_array</a>() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00099"></a><span class="lineno">   99</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a770f45838881fc061294e56d64f34386">num_arguments_per_1D_tensor</a>() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00104"></a><span class="lineno">  104</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac734502531e7f95e25b3bf688a304a59">num_arguments_per_2D_tensor</a>() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00109"></a><span class="lineno">  109</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a4feaae9c860cddfa843d37c953674a22">num_arguments_per_3D_tensor</a>() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00114"></a><span class="lineno">  114</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#af24df23a978e1a46ff6714e9b9d3eef9">num_arguments_per_4D_tensor</a>() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00122"></a><span class="lineno">  122</span>&#160;    <span class="keyword">virtual</span> <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#af6a174d47571f51f199ffc27ecc10f51">run</a>(<span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window, cl::CommandQueue &amp;queue) = 0;</div>
-<div class="line"><a name="l00128"></a><span class="lineno">  128</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
-<div class="line"><a name="l00129"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">  129</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">add_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, T <a class="code" href="hwc_8hpp.xhtml#a0f61d63b009d0880a89c843bd50d8d76">value</a>)</div>
-<div class="line"><a name="l00130"></a><span class="lineno">  130</span>&#160;    {</div>
-<div class="line"><a name="l00131"></a><span class="lineno">  131</span>&#160;        _kernel.setArg(idx++, value);</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;</div>
-<div class="line"><a name="l00140"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a44c701b9dbd01a171de4928254d1ecbf">  140</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a44c701b9dbd01a171de4928254d1ecbf">set_lws_hint</a>(cl::NDRange &amp;lws_hint)</div>
-<div class="line"><a name="l00141"></a><span class="lineno">  141</span>&#160;    {</div>
-<div class="line"><a name="l00142"></a><span class="lineno">  142</span>&#160;        _lws_hint = lws_hint;</div>
-<div class="line"><a name="l00143"></a><span class="lineno">  143</span>&#160;    }</div>
-<div class="line"><a name="l00144"></a><span class="lineno">  144</span>&#160;</div>
-<div class="line"><a name="l00155"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">  155</a></span>&#160;    <span class="keyword">const</span> std::string &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">config_id</a>()<span class="keyword"> const</span></div>
-<div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;<span class="keyword">    </span>{</div>
-<div class="line"><a name="l00157"></a><span class="lineno">  157</span>&#160;        <span class="keywordflow">return</span> _config_id;</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;</div>
-<div class="line"><a name="l00164"></a><span class="lineno">  164</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">set_target</a>(<a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a> target);</div>
-<div class="line"><a name="l00165"></a><span class="lineno">  165</span>&#160;</div>
-<div class="line"><a name="l00170"></a><span class="lineno">  170</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">set_target</a>(cl::Device &amp;device);</div>
-<div class="line"><a name="l00171"></a><span class="lineno">  171</span>&#160;</div>
-<div class="line"><a name="l00176"></a><span class="lineno">  176</span>&#160;    <a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">get_target</a>() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00177"></a><span class="lineno">  177</span>&#160;</div>
-<div class="line"><a name="l00178"></a><span class="lineno">  178</span>&#160;<span class="keyword">private</span>:</div>
-<div class="line"><a name="l00187"></a><span class="lineno">  187</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div>
-<div class="line"><a name="l00188"></a><span class="lineno">  188</span>&#160;    <span class="keywordtype">void</span> add_array_argument(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div>
-<div class="line"><a name="l00195"></a><span class="lineno">  195</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div>
-<div class="line"><a name="l00196"></a><span class="lineno">  196</span>&#160;    <span class="keywordtype">void</span> add_tensor_argument(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div>
-<div class="line"><a name="l00201"></a><span class="lineno">  201</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div>
-<div class="line"><a name="l00202"></a><span class="lineno">  202</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>           num_arguments_per_array() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00207"></a><span class="lineno">  207</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div>
-<div class="line"><a name="l00208"></a><span class="lineno">  208</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>           num_arguments_per_tensor() <span class="keyword">const</span>;</div>
-<div class="line"><a name="l00209"></a><span class="lineno">  209</span>&#160;</div>
-<div class="line"><a name="l00210"></a><span class="lineno">  210</span>&#160;<span class="keyword">protected</span>:</div>
-<div class="line"><a name="l00211"></a><span class="lineno">  211</span>&#160;    cl::Kernel  _kernel;    </div>
-<div class="line"><a name="l00212"></a><span class="lineno">  212</span>&#160;    cl::NDRange _lws_hint;  </div>
-<div class="line"><a name="l00213"></a><span class="lineno">  213</span>&#160;    <a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a>   _target;    </div>
-<div class="line"><a name="l00214"></a><span class="lineno">  214</span>&#160;    std::string _config_id; </div>
-<div class="line"><a name="l00215"></a><span class="lineno">  215</span>&#160;};</div>
-<div class="line"><a name="l00216"></a><span class="lineno">  216</span>&#160;</div>
-<div class="line"><a name="l00230"></a><span class="lineno">  230</span>&#160;<span class="keywordtype">void</span> <a class="code" href="namespacearm__compute.xhtml#a0f68015f649819f1102aa73d7b88b29e">enqueue</a>(cl::CommandQueue &amp;queue, ICLKernel &amp;kernel, <span class="keyword">const</span> Window &amp;window, <span class="keyword">const</span> cl::NDRange &amp;lws_hint = <a class="code" href="classarm__compute_1_1_c_l_kernel_library.xhtml#aebe2d25607545d31f2b1ace6d60a81da">CLKernelLibrary::get</a>().default_ndrange());</div>
-<div class="line"><a name="l00231"></a><span class="lineno">  231</span>&#160;</div>
-<div class="line"><a name="l00232"></a><span class="lineno">  232</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div>
-<div class="line"><a name="l00233"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a2d7c6b5f3332604ad6a637457f65c342">  233</a></span>&#160;<span class="keywordtype">void</span> ICLKernel::add_array_argument(<span class="keywordtype">unsigned</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window)</div>
-<div class="line"><a name="l00234"></a><span class="lineno">  234</span>&#160;{</div>
-<div class="line"><a name="l00235"></a><span class="lineno">  235</span>&#160;    <span class="comment">// Calculate offset to the start of the window</span></div>
-<div class="line"><a name="l00236"></a><span class="lineno">  236</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> offset_first_element = 0;</div>
-<div class="line"><a name="l00237"></a><span class="lineno">  237</span>&#160;</div>
-<div class="line"><a name="l00238"></a><span class="lineno">  238</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> n = 0; n &lt; num_dimensions; ++n)</div>
-<div class="line"><a name="l00239"></a><span class="lineno">  239</span>&#160;    {</div>
-<div class="line"><a name="l00240"></a><span class="lineno">  240</span>&#160;        offset_first_element += window[n].start() * strides[n];</div>
-<div class="line"><a name="l00241"></a><span class="lineno">  241</span>&#160;    }</div>
-<div class="line"><a name="l00242"></a><span class="lineno">  242</span>&#160;</div>
-<div class="line"><a name="l00243"></a><span class="lineno">  243</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> idx_start = idx;</div>
-<div class="line"><a name="l00244"></a><span class="lineno">  244</span>&#160;    _kernel.setArg(idx++, array-&gt;<a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml#a1fb4c50755a0ef424652246838ed91a6">cl_buffer</a>());</div>
-<div class="line"><a name="l00245"></a><span class="lineno">  245</span>&#160;</div>
-<div class="line"><a name="l00246"></a><span class="lineno">  246</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension = 0; dimension &lt; dimension_size; dimension++)</div>
-<div class="line"><a name="l00247"></a><span class="lineno">  247</span>&#160;    {</div>
-<div class="line"><a name="l00248"></a><span class="lineno">  248</span>&#160;        _kernel.setArg&lt;cl_uint&gt;(idx++, strides[dimension]);</div>
-<div class="line"><a name="l00249"></a><span class="lineno">  249</span>&#160;        _kernel.setArg&lt;cl_uint&gt;(idx++, strides[dimension] * window[dimension].step());</div>
-<div class="line"><a name="l00250"></a><span class="lineno">  250</span>&#160;    }</div>
-<div class="line"><a name="l00251"></a><span class="lineno">  251</span>&#160;</div>
-<div class="line"><a name="l00252"></a><span class="lineno">  252</span>&#160;    _kernel.setArg&lt;cl_uint&gt;(idx++, offset_first_element);</div>
-<div class="line"><a name="l00253"></a><span class="lineno">  253</span>&#160;</div>
-<div class="line"><a name="l00254"></a><span class="lineno">  254</span>&#160;    <a class="code" href="_error_8h.xhtml#a5bbdcf574d3f5e412fa6a1117911e67b">ARM_COMPUTE_ERROR_ON_MSG</a>(idx_start + num_arguments_per_array&lt;dimension_size&gt;() != idx,</div>
-<div class="line"><a name="l00255"></a><span class="lineno">  255</span>&#160;                             <span class="stringliteral">&quot;add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel&quot;</span>, dimension_size, num_arguments_per_array&lt;dimension_size&gt;());</div>
-<div class="line"><a name="l00256"></a><span class="lineno">  256</span>&#160;    <a class="code" href="_error_8h.xhtml#a4103adbb45806b2f2002d44b91d0d206">ARM_COMPUTE_UNUSED</a>(idx_start);</div>
-<div class="line"><a name="l00257"></a><span class="lineno">  257</span>&#160;}</div>
-<div class="line"><a name="l00258"></a><span class="lineno">  258</span>&#160;</div>
-<div class="line"><a name="l00259"></a><span class="lineno">  259</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
-<div class="line"><a name="l00260"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">  260</a></span>&#160;<span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">ICLKernel::add_1D_array_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window)</div>
-<div class="line"><a name="l00261"></a><span class="lineno">  261</span>&#160;{</div>
-<div class="line"><a name="l00262"></a><span class="lineno">  262</span>&#160;    add_array_argument&lt;T, 1&gt;(idx, array, strides, num_dimensions, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div>
-<div class="line"><a name="l00263"></a><span class="lineno">  263</span>&#160;}</div>
-<div class="line"><a name="l00264"></a><span class="lineno">  264</span>&#160;</div>
-<div class="line"><a name="l00265"></a><span class="lineno">  265</span>&#160;<span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div>
-<div class="line"><a name="l00266"></a><span class="lineno">  266</span>&#160;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>           ICLKernel::num_arguments_per_array()<span class="keyword"> const</span></div>
-<div class="line"><a name="l00267"></a><span class="lineno">  267</span>&#160;<span class="keyword"></span>{</div>
-<div class="line"><a name="l00268"></a><span class="lineno">  268</span>&#160;    <span class="keywordflow">return</span> num_arguments_per_tensor&lt;dimension_size&gt;();</div>
-<div class="line"><a name="l00269"></a><span class="lineno">  269</span>&#160;}</div>
-<div class="line"><a name="l00270"></a><span class="lineno">  270</span>&#160;</div>
-<div class="line"><a name="l00271"></a><span class="lineno">  271</span>&#160;<span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div>
-<div class="line"><a name="l00272"></a><span class="lineno">  272</span>&#160;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>           ICLKernel::num_arguments_per_tensor()<span class="keyword"> const</span></div>
-<div class="line"><a name="l00273"></a><span class="lineno">  273</span>&#160;<span class="keyword"></span>{</div>
-<div class="line"><a name="l00274"></a><span class="lineno">  274</span>&#160;    <span class="keywordflow">return</span> 2 + 2 * dimension_size;</div>
-<div class="line"><a name="l00275"></a><span class="lineno">  275</span>&#160;}</div>
-<div class="line"><a name="l00276"></a><span class="lineno">  276</span>&#160;}</div>
-<div class="line"><a name="l00277"></a><span class="lineno">  277</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/*__ARM_COMPUTE_ICLKERNEL_H__ */</span><span class="preprocessor"></span></div>
+<a href="_i_c_l_kernel_8h.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) 2016-2018 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">#ifndef __ARM_COMPUTE_ICLKERNEL_H__</span></div><div class="line"><a name="l00025"></a><span class="lineno">   25</span>&#160;<span class="preprocessor">#define __ARM_COMPUTE_ICLKERNEL_H__</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">#include &quot;<a class="code" href="_c_l_kernel_library_8h.xhtml">arm_compute/core/CL/CLKernelLibrary.h</a>&quot;</span></div><div class="line"><a name="l00028"></a><span class="lineno">   28</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_c_l_types_8h.xhtml">arm_compute/core/CL/CLTypes.h</a>&quot;</span></div><div class="line"><a name="l00029"></a><span class="lineno">   29</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_open_c_l_8h.xhtml">arm_compute/core/CL/OpenCL.h</a>&quot;</span></div><div class="line"><a name="l00030"></a><span class="lineno">   30</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_g_p_u_target_8h.xhtml">arm_compute/core/GPUTarget.h</a>&quot;</span></div><div class="line"><a name="l00031"></a><span class="lineno">   31</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_i_kernel_8h.xhtml">arm_compute/core/IKernel.h</a>&quot;</span></div><div class="line"><a name="l00032"></a><span class="lineno">   32</span>&#160;</div><div class="line"><a name="l00033"></a><span class="lineno">   33</span>&#160;<span class="preprocessor">#include &lt;string&gt;</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">   35</span>&#160;<span class="keyword">namespace </span><a class="code" href="namespacearm__compute.xhtml">arm_compute</a></div><div class="line"><a name="l00036"></a><span class="lineno">   36</span>&#160;{</div><div class="line"><a name="l00037"></a><span class="lineno">   37</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div><div class="line"><a name="l00038"></a><span class="lineno">   38</span>&#160;<span class="keyword">class </span>ICLArray;</div><div class="line"><a name="l00039"></a><span class="lineno">   39</span>&#160;<span class="keyword">class </span>ICLTensor;</div><div class="line"><a name="l00040"></a><span class="lineno">   40</span>&#160;<span class="keyword">class </span>Window;</div><div class="line"><a name="l00041"></a><span class="lineno">   41</span>&#160;</div><div class="line"><a name="l00043"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml">   43</a></span>&#160;<span class="keyword">class </span><a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml">ICLKernel</a> : <span class="keyword">public</span> <a class="code" href="classarm__compute_1_1_i_kernel.xhtml">IKernel</a></div><div class="line"><a name="l00044"></a><span class="lineno">   44</span>&#160;{</div><div class="line"><a name="l00045"></a><span class="lineno">   45</span>&#160;<span class="keyword">private</span>:</div><div class="line"><a name="l00050"></a><span class="lineno">   50</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>        dimension_size&gt;</div><div class="line"><a name="l00051"></a><span class="lineno">   51</span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_arguments_per_array()</div><div class="line"><a name="l00052"></a><span class="lineno">   52</span>&#160;    {</div><div class="line"><a name="l00053"></a><span class="lineno">   53</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;dimension_size&gt;();</div><div class="line"><a name="l00054"></a><span class="lineno">   54</span>&#160;    }</div><div class="line"><a name="l00059"></a><span class="lineno">   59</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>        dimension_size&gt;</div><div class="line"><a name="l00060"></a><span class="lineno">   60</span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_arguments_per_tensor()</div><div class="line"><a name="l00061"></a><span class="lineno">   61</span>&#160;    {</div><div class="line"><a name="l00062"></a><span class="lineno">   62</span>&#160;        <span class="keywordflow">return</span> 2 + 2 * dimension_size;</div><div class="line"><a name="l00063"></a><span class="lineno">   63</span>&#160;    }</div><div class="line"><a name="l00064"></a><span class="lineno">   64</span>&#160;</div><div class="line"><a name="l00065"></a><span class="lineno">   65</span>&#160;<span class="keyword">public</span>:</div><div class="line"><a name="l00067"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">   67</a></span>&#160;    <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">ICLKernel</a>()</div><div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;        : _kernel(nullptr), _lws_hint(<a class="code" href="classarm__compute_1_1_c_l_kernel_library.xhtml">CLKernelLibrary</a>::get().default_ndrange()), _target(<a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a>::<a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3afa362f4daec88442a387ff7cda411a38">MIDGARD</a>), _config_id(<a class="code" href="namespacearm__compute.xhtml">arm_compute</a>::default_config_id), _max_workgroup_size(0)</div><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;    }</div><div class="line"><a name="l00075"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">   75</a></span>&#160;    cl::Kernel &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">kernel</a>()</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="keywordflow">return</span> _kernel;</div><div class="line"><a name="l00078"></a><span class="lineno">   78</span>&#160;    }</div><div class="line"><a name="l00087"></a><span class="lineno">   87</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div><div class="line"><a name="l00088"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">   88</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">add_1D_array_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    {</div><div class="line"><a name="l00090"></a><span class="lineno">   90</span>&#160;        add_array_argument&lt;T, 1&gt;(idx, array, strides, num_dimensions, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00091"></a><span class="lineno">   91</span>&#160;    }</div><div class="line"><a name="l00098"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">   98</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">add_1D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</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;        add_tensor_argument&lt;1&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00101"></a><span class="lineno">  101</span>&#160;    }</div><div class="line"><a name="l00108"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">  108</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">add_2D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00109"></a><span class="lineno">  109</span>&#160;    {</div><div class="line"><a name="l00110"></a><span class="lineno">  110</span>&#160;        add_tensor_argument&lt;2&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00111"></a><span class="lineno">  111</span>&#160;    }</div><div class="line"><a name="l00118"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">  118</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">add_3D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00119"></a><span class="lineno">  119</span>&#160;    {</div><div class="line"><a name="l00120"></a><span class="lineno">  120</span>&#160;        add_tensor_argument&lt;3&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00121"></a><span class="lineno">  121</span>&#160;    }</div><div class="line"><a name="l00128"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">  128</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">add_4D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00129"></a><span class="lineno">  129</span>&#160;    {</div><div class="line"><a name="l00130"></a><span class="lineno">  130</span>&#160;        add_tensor_argument&lt;4&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00131"></a><span class="lineno">  131</span>&#160;    }</div><div class="line"><a name="l00136"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a278f0e6c68ca17e71f4c4ff82f360aa0">  136</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a278f0e6c68ca17e71f4c4ff82f360aa0">num_arguments_per_1D_array</a>()</div><div class="line"><a name="l00137"></a><span class="lineno">  137</span>&#160;    {</div><div class="line"><a name="l00138"></a><span class="lineno">  138</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_array&lt;1&gt;();</div><div class="line"><a name="l00139"></a><span class="lineno">  139</span>&#160;    }</div><div class="line"><a name="l00144"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a43b6c5e4b57069c5f61e96dff24c212d">  144</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a43b6c5e4b57069c5f61e96dff24c212d">num_arguments_per_1D_tensor</a>()</div><div class="line"><a name="l00145"></a><span class="lineno">  145</span>&#160;    {</div><div class="line"><a name="l00146"></a><span class="lineno">  146</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;1&gt;();</div><div class="line"><a name="l00147"></a><span class="lineno">  147</span>&#160;    }</div><div class="line"><a name="l00152"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a45601e0d46621a5b6f2e417d60e5c800">  152</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a45601e0d46621a5b6f2e417d60e5c800">num_arguments_per_2D_tensor</a>()</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;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;2&gt;();</div><div class="line"><a name="l00155"></a><span class="lineno">  155</span>&#160;    }</div><div class="line"><a name="l00160"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6c9c1e7a7d96743375ca40847f0f12e2">  160</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6c9c1e7a7d96743375ca40847f0f12e2">num_arguments_per_3D_tensor</a>()</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;    {</div><div class="line"><a name="l00162"></a><span class="lineno">  162</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;3&gt;();</div><div class="line"><a name="l00163"></a><span class="lineno">  163</span>&#160;    }</div><div class="line"><a name="l00168"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a184fdf37587a9314cf12623accea6c73">  168</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a184fdf37587a9314cf12623accea6c73">num_arguments_per_4D_tensor</a>()</div><div class="line"><a name="l00169"></a><span class="lineno">  169</span>&#160;    {</div><div class="line"><a name="l00170"></a><span class="lineno">  170</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;4&gt;();</div><div class="line"><a name="l00171"></a><span class="lineno">  171</span>&#160;    }</div><div class="line"><a name="l00179"></a><span class="lineno">  179</span>&#160;    <span class="keyword">virtual</span> <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#af6a174d47571f51f199ffc27ecc10f51">run</a>(<span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>, cl::CommandQueue &amp;queue) = 0;</div><div class="line"><a name="l00185"></a><span class="lineno">  185</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div><div class="line"><a name="l00186"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">  186</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">add_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, T value)</div><div class="line"><a name="l00187"></a><span class="lineno">  187</span>&#160;    {</div><div class="line"><a name="l00188"></a><span class="lineno">  188</span>&#160;        _kernel.setArg(idx++, value);</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;</div><div class="line"><a name="l00197"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad356b88c8f61267d593d9ed99835bde9">  197</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad356b88c8f61267d593d9ed99835bde9">set_lws_hint</a>(<span class="keyword">const</span> cl::NDRange &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a>)</div><div class="line"><a name="l00198"></a><span class="lineno">  198</span>&#160;    {</div><div class="line"><a name="l00199"></a><span class="lineno">  199</span>&#160;        _lws_hint = <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a>;</div><div class="line"><a name="l00200"></a><span class="lineno">  200</span>&#160;    }</div><div class="line"><a name="l00201"></a><span class="lineno">  201</span>&#160;</div><div class="line"><a name="l00206"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">  206</a></span>&#160;    cl::NDRange <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a>()<span class="keyword"> const</span></div><div class="line"><a name="l00207"></a><span class="lineno">  207</span>&#160;<span class="keyword">    </span>{</div><div class="line"><a name="l00208"></a><span class="lineno">  208</span>&#160;        <span class="keywordflow">return</span> _lws_hint;</div><div class="line"><a name="l00209"></a><span class="lineno">  209</span>&#160;    }</div><div class="line"><a name="l00210"></a><span class="lineno">  210</span>&#160;</div><div class="line"><a name="l00221"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">  221</a></span>&#160;    <span class="keyword">const</span> std::string &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">config_id</a>()<span class="keyword"> const</span></div><div class="line"><a name="l00222"></a><span class="lineno">  222</span>&#160;<span class="keyword">    </span>{</div><div class="line"><a name="l00223"></a><span class="lineno">  223</span>&#160;        <span class="keywordflow">return</span> _config_id;</div><div class="line"><a name="l00224"></a><span class="lineno">  224</span>&#160;    }</div><div class="line"><a name="l00225"></a><span class="lineno">  225</span>&#160;</div><div class="line"><a name="l00230"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">  230</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">set_target</a>(<a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a> target)</div><div class="line"><a name="l00231"></a><span class="lineno">  231</span>&#160;    {</div><div class="line"><a name="l00232"></a><span class="lineno">  232</span>&#160;        _target = target;</div><div class="line"><a name="l00233"></a><span class="lineno">  233</span>&#160;    }</div><div class="line"><a name="l00234"></a><span class="lineno">  234</span>&#160;</div><div class="line"><a name="l00239"></a><span class="lineno">  239</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">set_target</a>(cl::Device &amp;device);</div><div class="line"><a name="l00240"></a><span class="lineno">  240</span>&#160;</div><div class="line"><a name="l00245"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">  245</a></span>&#160;    <a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">get_target</a>()<span class="keyword"> const</span></div><div class="line"><a name="l00246"></a><span class="lineno">  246</span>&#160;<span class="keyword">    </span>{</div><div class="line"><a name="l00247"></a><span class="lineno">  247</span>&#160;        <span class="keywordflow">return</span> _target;</div><div class="line"><a name="l00248"></a><span class="lineno">  248</span>&#160;    }</div><div class="line"><a name="l00249"></a><span class="lineno">  249</span>&#160;</div><div class="line"><a name="l00254"></a><span class="lineno">  254</span>&#160;    <span class="keywordtype">size_t</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#abca336f832d730e8494049bd714df60a">get_max_workgroup_size</a>();</div><div class="line"><a name="l00261"></a><span class="lineno">  261</span>&#160;    <span class="keyword">static</span> cl::NDRange <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a4b4a20e518c04e657839280cd486a85d">gws_from_window</a>(<span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div><div class="line"><a name="l00262"></a><span class="lineno">  262</span>&#160;</div><div class="line"><a name="l00263"></a><span class="lineno">  263</span>&#160;<span class="keyword">private</span>:</div><div class="line"><a name="l00272"></a><span class="lineno">  272</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div><div class="line"><a name="l00273"></a><span class="lineno">  273</span>&#160;    <span class="keywordtype">void</span> add_array_argument(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div><div class="line"><a name="l00280"></a><span class="lineno">  280</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div><div class="line"><a name="l00281"></a><span class="lineno">  281</span>&#160;    <span class="keywordtype">void</span> add_tensor_argument(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div><div class="line"><a name="l00282"></a><span class="lineno">  282</span>&#160;</div><div class="line"><a name="l00283"></a><span class="lineno">  283</span>&#160;<span class="keyword">protected</span>:</div><div class="line"><a name="l00284"></a><span class="lineno">  284</span>&#160;    cl::Kernel  _kernel;             </div><div class="line"><a name="l00285"></a><span class="lineno">  285</span>&#160;    cl::NDRange _lws_hint;           </div><div class="line"><a name="l00286"></a><span class="lineno">  286</span>&#160;    <a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a>   _target;             </div><div class="line"><a name="l00287"></a><span class="lineno">  287</span>&#160;    std::string _config_id;          </div><div class="line"><a name="l00288"></a><span class="lineno">  288</span>&#160;    <span class="keywordtype">size_t</span>      _max_workgroup_size; </div><div class="line"><a name="l00289"></a><span class="lineno">  289</span>&#160;};</div><div class="line"><a name="l00290"></a><span class="lineno">  290</span>&#160;</div><div class="line"><a name="l00304"></a><span class="lineno">  304</span>&#160;<span class="keywordtype">void</span> <a class="code" href="namespacearm__compute.xhtml#a0f68015f649819f1102aa73d7b88b29e">enqueue</a>(cl::CommandQueue &amp;queue, <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml">ICLKernel</a> &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">kernel</a>, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>, <span class="keyword">const</span> cl::NDRange &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a> = <a class="code" href="classarm__compute_1_1_c_l_kernel_library.xhtml#aebe2d25607545d31f2b1ace6d60a81da">CLKernelLibrary::get</a>().default_ndrange());</div><div class="line"><a name="l00305"></a><span class="lineno">  305</span>&#160;</div><div class="line"><a name="l00314"></a><span class="lineno">  314</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div><div class="line"><a name="l00315"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a2d7c6b5f3332604ad6a637457f65c342">  315</a></span>&#160;<span class="keywordtype">void</span> ICLKernel::add_array_argument(<span class="keywordtype">unsigned</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window)</div><div class="line"><a name="l00316"></a><span class="lineno">  316</span>&#160;{</div><div class="line"><a name="l00317"></a><span class="lineno">  317</span>&#160;    <a class="code" href="_error_8h.xhtml#a54a6080c9f4df1f908e57a9bbb46f5da">ARM_COMPUTE_ERROR_ON</a>(array == <span class="keyword">nullptr</span>);</div><div class="line"><a name="l00318"></a><span class="lineno">  318</span>&#160;</div><div class="line"><a name="l00319"></a><span class="lineno">  319</span>&#160;    <span class="comment">// Calculate offset to the start of the window</span></div><div class="line"><a name="l00320"></a><span class="lineno">  320</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> offset_first_element = 0;</div><div class="line"><a name="l00321"></a><span class="lineno">  321</span>&#160;</div><div class="line"><a name="l00322"></a><span class="lineno">  322</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> n = 0; n &lt; num_dimensions; ++n)</div><div class="line"><a name="l00323"></a><span class="lineno">  323</span>&#160;    {</div><div class="line"><a name="l00324"></a><span class="lineno">  324</span>&#160;        offset_first_element += window[n].start() * strides[n];</div><div class="line"><a name="l00325"></a><span class="lineno">  325</span>&#160;    }</div><div class="line"><a name="l00326"></a><span class="lineno">  326</span>&#160;</div><div class="line"><a name="l00327"></a><span class="lineno">  327</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> idx_start = idx;</div><div class="line"><a name="l00328"></a><span class="lineno">  328</span>&#160;    _kernel.setArg(idx++, array-&gt;<a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml#a1fb4c50755a0ef424652246838ed91a6">cl_buffer</a>());</div><div class="line"><a name="l00329"></a><span class="lineno">  329</span>&#160;</div><div class="line"><a name="l00330"></a><span class="lineno">  330</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension = 0; dimension &lt; dimension_size; dimension++)</div><div class="line"><a name="l00331"></a><span class="lineno">  331</span>&#160;    {</div><div class="line"><a name="l00332"></a><span class="lineno">  332</span>&#160;        _kernel.setArg&lt;cl_uint&gt;(idx++, strides[dimension]);</div><div class="line"><a name="l00333"></a><span class="lineno">  333</span>&#160;        _kernel.setArg&lt;cl_uint&gt;(idx++, strides[dimension] * window[dimension].step());</div><div class="line"><a name="l00334"></a><span class="lineno">  334</span>&#160;    }</div><div class="line"><a name="l00335"></a><span class="lineno">  335</span>&#160;</div><div class="line"><a name="l00336"></a><span class="lineno">  336</span>&#160;    _kernel.setArg&lt;cl_uint&gt;(idx++, offset_first_element);</div><div class="line"><a name="l00337"></a><span class="lineno">  337</span>&#160;</div><div class="line"><a name="l00338"></a><span class="lineno">  338</span>&#160;    <a class="code" href="_error_8h.xhtml#a5bbdcf574d3f5e412fa6a1117911e67b">ARM_COMPUTE_ERROR_ON_MSG</a>(idx_start + num_arguments_per_array&lt;dimension_size&gt;() != idx,</div><div class="line"><a name="l00339"></a><span class="lineno">  339</span>&#160;                             <span class="stringliteral">&quot;add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel&quot;</span>, dimension_size, num_arguments_per_array&lt;dimension_size&gt;());</div><div class="line"><a name="l00340"></a><span class="lineno">  340</span>&#160;    <a class="code" href="_error_8h.xhtml#a6dc630a6ae9cc063b3924bcea8dee9d6">ARM_COMPUTE_UNUSED</a>(idx_start);</div><div class="line"><a name="l00341"></a><span class="lineno">  341</span>&#160;}</div><div class="line"><a name="l00342"></a><span class="lineno">  342</span>&#160;}</div><div class="line"><a name="l00343"></a><span class="lineno">  343</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/*__ARM_COMPUTE_ICLKERNEL_H__ */</span><span class="preprocessor"></span></div><div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a43b6c5e4b57069c5f61e96dff24c212d"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a43b6c5e4b57069c5f61e96dff24c212d">arm_compute::ICLKernel::num_arguments_per_1D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_1D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 1D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00144">ICLKernel.h:144</a></div></div>
 <div class="ttc" id="classarm__compute_1_1_i_kernel_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_kernel.xhtml">arm_compute::IKernel</a></div><div class="ttdoc">Common information for all the kernels. </div><div class="ttdef"><b>Definition:</b> <a href="_i_kernel_8h_source.xhtml#l00033">IKernel.h:33</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a4feaae9c860cddfa843d37c953674a22"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a4feaae9c860cddfa843d37c953674a22">arm_compute::ICLKernel::num_arguments_per_3D_tensor</a></div><div class="ttdeci">unsigned int num_arguments_per_3D_tensor() const </div><div class="ttdoc">Returns the number of arguments enqueued per 3D tensor object. </div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ae5121015ab09ece4d470f50c7ffe198e"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">arm_compute::ICLKernel::kernel</a></div><div class="ttdeci">cl::Kernel &amp; kernel()</div><div class="ttdoc">Returns a reference to the OpenCL kernel of this object. </div></div>
-<div class="ttc" id="_error_8h_xhtml_a4103adbb45806b2f2002d44b91d0d206"><div class="ttname"><a href="_error_8h.xhtml#a4103adbb45806b2f2002d44b91d0d206">ARM_COMPUTE_UNUSED</a></div><div class="ttdeci">#define ARM_COMPUTE_UNUSED(var)</div><div class="ttdoc">To avoid unused variables warnings. </div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00049">Error.h:49</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a50f427a1d9419800972b9e03c4034311"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">arm_compute::ICLKernel::add_argument</a></div><div class="ttdeci">void add_argument(unsigned int &amp;idx, T value)</div><div class="ttdoc">Add the passed parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00129">ICLKernel.h:129</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a278f0e6c68ca17e71f4c4ff82f360aa0"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a278f0e6c68ca17e71f4c4ff82f360aa0">arm_compute::ICLKernel::num_arguments_per_1D_array</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_1D_array()</div><div class="ttdoc">Returns the number of arguments enqueued per 1D array object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00136">ICLKernel.h:136</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a0737b867334392d5dc5c2e3291d67dbf"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">arm_compute::ICLKernel::lws_hint</a></div><div class="ttdeci">cl::NDRange lws_hint() const </div><div class="ttdoc">Return the Local-Workgroup-Size hint. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00206">ICLKernel.h:206</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ae5121015ab09ece4d470f50c7ffe198e"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">arm_compute::ICLKernel::kernel</a></div><div class="ttdeci">cl::Kernel &amp; kernel()</div><div class="ttdoc">Returns a reference to the OpenCL kernel of this object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00075">ICLKernel.h:75</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ad356b88c8f61267d593d9ed99835bde9"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad356b88c8f61267d593d9ed99835bde9">arm_compute::ICLKernel::set_lws_hint</a></div><div class="ttdeci">void set_lws_hint(const cl::NDRange &amp;lws_hint)</div><div class="ttdoc">Set the Local-Workgroup-Size hint. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00197">ICLKernel.h:197</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a50f427a1d9419800972b9e03c4034311"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">arm_compute::ICLKernel::add_argument</a></div><div class="ttdeci">void add_argument(unsigned int &amp;idx, T value)</div><div class="ttdoc">Add the passed parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00186">ICLKernel.h:186</a></div></div>
 <div class="ttc" id="_c_l_types_8h_xhtml"><div class="ttname"><a href="_c_l_types_8h.xhtml">CLTypes.h</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a9331d385192a50adf74d3af40ce0fa20"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">arm_compute::ICLKernel::add_1D_array_argument</a></div><div class="ttdeci">void add_1D_array_argument(unsigned int &amp;idx, const ICLArray&lt; T &gt; *array, const Strides &amp;strides, unsigned int num_dimensions, const Window &amp;window)</div><div class="ttdoc">Add the passed 1D array&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00260">ICLKernel.h:260</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a770f45838881fc061294e56d64f34386"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a770f45838881fc061294e56d64f34386">arm_compute::ICLKernel::num_arguments_per_1D_tensor</a></div><div class="ttdeci">unsigned int num_arguments_per_1D_tensor() const </div><div class="ttdoc">Returns the number of arguments enqueued per 1D tensor object. </div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a9331d385192a50adf74d3af40ce0fa20"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">arm_compute::ICLKernel::add_1D_array_argument</a></div><div class="ttdeci">void add_1D_array_argument(unsigned int &amp;idx, const ICLArray&lt; T &gt; *array, const Strides &amp;strides, unsigned int num_dimensions, const Window &amp;window)</div><div class="ttdoc">Add the passed 1D array&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00088">ICLKernel.h:88</a></div></div>
+<div class="ttc" id="_error_8h_xhtml_a54a6080c9f4df1f908e57a9bbb46f5da"><div class="ttname"><a href="_error_8h.xhtml#a54a6080c9f4df1f908e57a9bbb46f5da">ARM_COMPUTE_ERROR_ON</a></div><div class="ttdeci">#define ARM_COMPUTE_ERROR_ON(cond)</div><div class="ttdoc">If the condition is true then an error message is printed and an exception thrown. </div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00328">Error.h:328</a></div></div>
 <div class="ttc" id="_i_kernel_8h_xhtml"><div class="ttname"><a href="_i_kernel_8h.xhtml">IKernel.h</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml">arm_compute::ICLKernel</a></div><div class="ttdoc">Common interface for all the OpenCL kernels. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00042">ICLKernel.h:42</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a28f5847162f352444c6ac1825d0e99c7"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">arm_compute::ICLKernel::add_3D_tensor_argument</a></div><div class="ttdeci">void add_3D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 3D tensor&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...</div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml">arm_compute::ICLKernel</a></div><div class="ttdoc">Common interface for all the OpenCL kernels. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00043">ICLKernel.h:43</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a28f5847162f352444c6ac1825d0e99c7"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">arm_compute::ICLKernel::add_3D_tensor_argument</a></div><div class="ttdeci">void add_3D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 3D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00118">ICLKernel.h:118</a></div></div>
+<div class="ttc" id="namespacearm__compute_xhtml"><div class="ttname"><a href="namespacearm__compute.xhtml">arm_compute</a></div><div class="ttdoc">This file contains all available output stages for GEMMLowp on OpenCL. </div><div class="ttdef"><b>Definition:</b> <a href="00__introduction_8dox_source.xhtml#l00001">00_introduction.dox:1</a></div></div>
 <div class="ttc" id="classarm__compute_1_1_c_l_kernel_library_xhtml_aebe2d25607545d31f2b1ace6d60a81da"><div class="ttname"><a href="classarm__compute_1_1_c_l_kernel_library.xhtml#aebe2d25607545d31f2b1ace6d60a81da">arm_compute::CLKernelLibrary::get</a></div><div class="ttdeci">static CLKernelLibrary &amp; get()</div><div class="ttdoc">Access the KernelLibrary singleton. </div></div>
 <div class="ttc" id="classarm__compute_1_1_i_c_l_array_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_array.xhtml">arm_compute::ICLArray</a></div><div class="ttdoc">Interface for OpenCL Array. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_array_8h_source.xhtml#l00035">ICLArray.h:35</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ac734502531e7f95e25b3bf688a304a59"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac734502531e7f95e25b3bf688a304a59">arm_compute::ICLKernel::num_arguments_per_2D_tensor</a></div><div class="ttdeci">unsigned int num_arguments_per_2D_tensor() const </div><div class="ttdoc">Returns the number of arguments enqueued per 2D tensor object. </div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a44c701b9dbd01a171de4928254d1ecbf"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a44c701b9dbd01a171de4928254d1ecbf">arm_compute::ICLKernel::set_lws_hint</a></div><div class="ttdeci">void set_lws_hint(cl::NDRange &amp;lws_hint)</div><div class="ttdoc">Set the Local-Workgroup-Size hint. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00140">ICLKernel.h:140</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a646cd535a16835b246c3367a63d96250"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">arm_compute::ICLKernel::get_target</a></div><div class="ttdeci">GPUTarget get_target() const </div><div class="ttdoc">Get the targeted GPU architecture. </div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a226d2e9e6d3c42d681666566fe950b2d"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">arm_compute::ICLKernel::config_id</a></div><div class="ttdeci">const std::string &amp; config_id() const </div><div class="ttdoc">Get the configuration ID. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00155">ICLKernel.h:155</a></div></div>
-<div class="ttc" id="hwc_8hpp_xhtml_a0f61d63b009d0880a89c843bd50d8d76"><div class="ttname"><a href="hwc_8hpp.xhtml#a0f61d63b009d0880a89c843bd50d8d76">value</a></div><div class="ttdeci">void * value</div><div class="ttdef"><b>Definition:</b> <a href="hwc_8hpp_source.xhtml#l00269">hwc.hpp:269</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_strides_xhtml"><div class="ttname"><a href="classarm__compute_1_1_strides.xhtml">arm_compute::Strides</a></div><div class="ttdoc">Strides of an item in bytes. </div><div class="ttdef"><b>Definition:</b> <a href="_strides_8h_source.xhtml#l00038">Strides.h:38</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a6c9c1e7a7d96743375ca40847f0f12e2"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6c9c1e7a7d96743375ca40847f0f12e2">arm_compute::ICLKernel::num_arguments_per_3D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_3D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 3D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00160">ICLKernel.h:160</a></div></div>
+<div class="ttc" id="_error_8h_xhtml_a6dc630a6ae9cc063b3924bcea8dee9d6"><div class="ttname"><a href="_error_8h.xhtml#a6dc630a6ae9cc063b3924bcea8dee9d6">ARM_COMPUTE_UNUSED</a></div><div class="ttdeci">#define ARM_COMPUTE_UNUSED(...)</div><div class="ttdoc">To avoid unused variables warnings. </div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00159">Error.h:159</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a646cd535a16835b246c3367a63d96250"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">arm_compute::ICLKernel::get_target</a></div><div class="ttdeci">GPUTarget get_target() const </div><div class="ttdoc">Get the targeted GPU architecture. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00245">ICLKernel.h:245</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a45601e0d46621a5b6f2e417d60e5c800"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a45601e0d46621a5b6f2e417d60e5c800">arm_compute::ICLKernel::num_arguments_per_2D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_2D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 2D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00152">ICLKernel.h:152</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a184fdf37587a9314cf12623accea6c73"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a184fdf37587a9314cf12623accea6c73">arm_compute::ICLKernel::num_arguments_per_4D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_4D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 4D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00168">ICLKernel.h:168</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a226d2e9e6d3c42d681666566fe950b2d"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">arm_compute::ICLKernel::config_id</a></div><div class="ttdeci">const std::string &amp; config_id() const </div><div class="ttdoc">Get the configuration ID. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00221">ICLKernel.h:221</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a4b4a20e518c04e657839280cd486a85d"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a4b4a20e518c04e657839280cd486a85d">arm_compute::ICLKernel::gws_from_window</a></div><div class="ttdeci">static cl::NDRange gws_from_window(const Window &amp;window)</div><div class="ttdoc">Get the global work size given an execution window. </div></div>
+<div class="ttc" id="classarm__compute_1_1_strides_xhtml"><div class="ttname"><a href="classarm__compute_1_1_strides.xhtml">arm_compute::Strides</a></div><div class="ttdoc">Strides of an item in bytes. </div><div class="ttdef"><b>Definition:</b> <a href="_strides_8h_source.xhtml#l00037">Strides.h:37</a></div></div>
 <div class="ttc" id="namespacearm__compute_xhtml_a0f68015f649819f1102aa73d7b88b29e"><div class="ttname"><a href="namespacearm__compute.xhtml#a0f68015f649819f1102aa73d7b88b29e">arm_compute::enqueue</a></div><div class="ttdeci">void enqueue(cl::CommandQueue &amp;queue, ICLKernel &amp;kernel, const Window &amp;window, const cl::NDRange &amp;lws_hint=CLKernelLibrary::get().default_ndrange())</div><div class="ttdoc">Add the kernel to the command queue with the given window. </div></div>
-<div class="ttc" id="namespacearm__compute_xhtml_af87af588094f04e537df98eaf0426cb6"><div class="ttname"><a href="namespacearm__compute.xhtml#af87af588094f04e537df98eaf0426cb6">arm_compute::ICLTensor</a></div><div class="ttdeci">ICLImage ICLTensor</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_histogram_8h_source.xhtml#l00034">CLHistogram.h:34</a></div></div>
+<div class="ttc" id="_g_p_u_target_8h_xhtml"><div class="ttname"><a href="_g_p_u_target_8h.xhtml">GPUTarget.h</a></div></div>
 <div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_af6a174d47571f51f199ffc27ecc10f51"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#af6a174d47571f51f199ffc27ecc10f51">arm_compute::ICLKernel::run</a></div><div class="ttdeci">virtual void run(const Window &amp;window, cl::CommandQueue &amp;queue)=0</div><div class="ttdoc">Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue...</div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ac74dad3e61f79334f5e73f3c3ac603cb"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">arm_compute::ICLKernel::add_2D_tensor_argument</a></div><div class="ttdeci">void add_2D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 2D tensor&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...</div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ac74dad3e61f79334f5e73f3c3ac603cb"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">arm_compute::ICLKernel::add_2D_tensor_argument</a></div><div class="ttdeci">void add_2D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 2D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00108">ICLKernel.h:108</a></div></div>
 <div class="ttc" id="classarm__compute_1_1_i_c_l_tensor_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_tensor.xhtml">arm_compute::ICLTensor</a></div><div class="ttdoc">Interface for OpenCL tensor. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_tensor_8h_source.xhtml#l00040">ICLTensor.h:40</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a6b10e96ce90bf901d17def86b874b019"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">arm_compute::ICLKernel::ICLKernel</a></div><div class="ttdeci">ICLKernel()</div><div class="ttdoc">Constructor. </div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a99fde125501fae43952af222c91236cd"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a99fde125501fae43952af222c91236cd">arm_compute::ICLKernel::num_arguments_per_1D_array</a></div><div class="ttdeci">unsigned int num_arguments_per_1D_array() const </div><div class="ttdoc">Returns the number of arguments enqueued per 1D array object. </div></div>
-<div class="ttc" id="namespacearm__compute_xhtml_a735ac6c2a02e320969625308810444f3"><div class="ttname"><a href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">arm_compute::GPUTarget</a></div><div class="ttdeci">GPUTarget</div><div class="ttdoc">Available GPU Targets. </div><div class="ttdef"><b>Definition:</b> <a href="_c_l_types_8h_source.xhtml#l00035">CLTypes.h:35</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a6b10e96ce90bf901d17def86b874b019"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">arm_compute::ICLKernel::ICLKernel</a></div><div class="ttdeci">ICLKernel()</div><div class="ttdoc">Constructor. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00067">ICLKernel.h:67</a></div></div>
+<div class="ttc" id="namespacearm__compute_xhtml_a735ac6c2a02e320969625308810444f3"><div class="ttname"><a href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">arm_compute::GPUTarget</a></div><div class="ttdeci">GPUTarget</div><div class="ttdoc">Available GPU Targets. </div><div class="ttdef"><b>Definition:</b> <a href="_g_p_u_target_8h_source.xhtml#l00034">GPUTarget.h:34</a></div></div>
 <div class="ttc" id="_c_l_kernel_library_8h_xhtml"><div class="ttname"><a href="_c_l_kernel_library_8h.xhtml">CLKernelLibrary.h</a></div></div>
+<div class="ttc" id="namespacearm__compute_xhtml_a735ac6c2a02e320969625308810444f3afa362f4daec88442a387ff7cda411a38"><div class="ttname"><a href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3afa362f4daec88442a387ff7cda411a38">arm_compute::GPUTarget::MIDGARD</a></div></div>
 <div class="ttc" id="_open_c_l_8h_xhtml"><div class="ttname"><a href="_open_c_l_8h.xhtml">OpenCL.h</a></div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ad5ba9d34a3a855bf1dd2e36316ff550a"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">arm_compute::ICLKernel::set_target</a></div><div class="ttdeci">void set_target(GPUTarget target)</div><div class="ttdoc">Set the targeted GPU architecture. </div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a479e7043e65dc87de35d374e108510f7"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">arm_compute::ICLKernel::add_1D_tensor_argument</a></div><div class="ttdeci">void add_1D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 1D tensor&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...</div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a33e09c946b338fbfc780a9d1c66e68ad"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">arm_compute::ICLKernel::add_4D_tensor_argument</a></div><div class="ttdeci">void add_4D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 4D tensor&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...</div></div>
+<div class="ttc" id="classarm__compute_1_1_c_l_kernel_library_xhtml"><div class="ttname"><a href="classarm__compute_1_1_c_l_kernel_library.xhtml">arm_compute::CLKernelLibrary</a></div><div class="ttdoc">CLKernelLibrary class. </div><div class="ttdef"><b>Definition:</b> <a href="_c_l_kernel_library_8h_source.xhtml#l00192">CLKernelLibrary.h:192</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_abca336f832d730e8494049bd714df60a"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#abca336f832d730e8494049bd714df60a">arm_compute::ICLKernel::get_max_workgroup_size</a></div><div class="ttdeci">size_t get_max_workgroup_size()</div><div class="ttdoc">Get the maximum workgroup size for the device the CLKernelLibrary uses. </div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ad5ba9d34a3a855bf1dd2e36316ff550a"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">arm_compute::ICLKernel::set_target</a></div><div class="ttdeci">void set_target(GPUTarget target)</div><div class="ttdoc">Set the targeted GPU architecture. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00230">ICLKernel.h:230</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a479e7043e65dc87de35d374e108510f7"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">arm_compute::ICLKernel::add_1D_tensor_argument</a></div><div class="ttdeci">void add_1D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 1D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00098">ICLKernel.h:98</a></div></div>
+<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a33e09c946b338fbfc780a9d1c66e68ad"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">arm_compute::ICLKernel::add_4D_tensor_argument</a></div><div class="ttdeci">void add_4D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 4D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00128">ICLKernel.h:128</a></div></div>
 <div class="ttc" id="classarm__compute_1_1_i_c_l_array_xhtml_a1fb4c50755a0ef424652246838ed91a6"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_array.xhtml#a1fb4c50755a0ef424652246838ed91a6">arm_compute::ICLArray::cl_buffer</a></div><div class="ttdeci">virtual const cl::Buffer &amp; cl_buffer() const =0</div><div class="ttdoc">Interface to be implemented by the child class to return a reference to the OpenCL buffer containing ...</div></div>
-<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_af24df23a978e1a46ff6714e9b9d3eef9"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#af24df23a978e1a46ff6714e9b9d3eef9">arm_compute::ICLKernel::num_arguments_per_4D_tensor</a></div><div class="ttdeci">unsigned int num_arguments_per_4D_tensor() const </div><div class="ttdoc">Returns the number of arguments enqueued per 4D tensor object. </div></div>
 <div class="ttc" id="classarm__compute_1_1_i_kernel_xhtml_a3f5646133956f06348b310ccc3d36353"><div class="ttname"><a href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">arm_compute::IKernel::window</a></div><div class="ttdeci">const Window &amp; window() const </div><div class="ttdoc">The maximum window the kernel can be executed on. </div></div>
 <div class="ttc" id="classarm__compute_1_1_window_xhtml"><div class="ttname"><a href="classarm__compute_1_1_window.xhtml">arm_compute::Window</a></div><div class="ttdoc">Describe a multidimensional execution window. </div><div class="ttdef"><b>Definition:</b> <a href="_window_8h_source.xhtml#l00039">Window.h:39</a></div></div>
-<div class="ttc" id="_error_8h_xhtml_a5bbdcf574d3f5e412fa6a1117911e67b"><div class="ttname"><a href="_error_8h.xhtml#a5bbdcf574d3f5e412fa6a1117911e67b">ARM_COMPUTE_ERROR_ON_MSG</a></div><div class="ttdeci">#define ARM_COMPUTE_ERROR_ON_MSG(cond,...)</div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00115">Error.h:115</a></div></div>
+<div class="ttc" id="_error_8h_xhtml_a5bbdcf574d3f5e412fa6a1117911e67b"><div class="ttname"><a href="_error_8h.xhtml#a5bbdcf574d3f5e412fa6a1117911e67b">ARM_COMPUTE_ERROR_ON_MSG</a></div><div class="ttdeci">#define ARM_COMPUTE_ERROR_ON_MSG(cond,...)</div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00319">Error.h:319</a></div></div>
 </div><!-- fragment --></div><!-- contents -->
 </div><!-- doc-content -->
 <!-- start footer part -->
 <div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
   <ul>
-    <li class="navelem"><a class="el" href="dir_d4d363c0979c7ffda4833df3f93b5668.xhtml">arm_compute</a></li><li class="navelem"><a class="el" href="dir_62e76289bb143101bef0ade0babfd908.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_05f0d9b6fc8ae7732f7cdd827386fe59.xhtml">CL</a></li><li class="navelem"><a class="el" href="_i_c_l_kernel_8h.xhtml">ICLKernel.h</a></li>
-    <li class="footer">Generated on Thu Sep 28 2017 14:37:49 for Compute Library by
+    <li class="navelem"><a class="el" href="dir_214608ef36d61442cb2b0c1c4e9a7def.xhtml">arm_compute</a></li><li class="navelem"><a class="el" href="dir_1fb090f0c6070330bfaccc4236d3ca0d.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_1633709b037eda5b8477677909f78129.xhtml">CL</a></li><li class="navelem"><a class="el" href="_i_c_l_kernel_8h.xhtml">ICLKernel.h</a></li>
+    <li class="footer">Generated on Wed May 23 2018 11:36:34 for Compute Library by
     <a href="http://www.doxygen.org/index.html">
-    <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.6 </li>
+    <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
   </ul>
 </div>
 </body>