documentation/pooling__layer_8cl.xhtml - platform/external/ComputeLibrary - Git at Google

 <!-- HTML header for doxygen 1.8.15-->
 <!-- Remember to use version doxygen 1.8.15 +-->
 <!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
 <html xmlns="http://www.w3.org/1999/xhtml">
 <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.15"/>
 <meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
 <title>Compute Library: src/core/CL/cl_kernels/pooling_layer.cl File Reference</title>
 <link href="tabs.css" rel="stylesheet" type="text/css"/>
 <script type="text/javascript" src="jquery.js"></script>
 <script type="text/javascript" src="dynsections.js"></script>
 <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">
 /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
   $(document).ready(initResizable);
 /* @license-end */</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/x-mathjax-config">
   MathJax.Hub.Config({
     extensions: ["tex2jax.js"],
     jax: ["input/TeX","output/HTML-CSS"],
 });
 </script><script type="text/javascript" async="async" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
 <link href="doxygen.css" rel="stylesheet" type="text/css" />
 <link href="stylesheet.css" rel="stylesheet" type="text/css"/>
 </head>
 <body>
 <div id="top"><!-- do not remove this div, it is closed by doxygen! -->
 <div id="titlearea">
 <table cellspacing="0" cellpadding="0">
  <tbody>
  <tr style="height: 56px;">
   <img alt="Compute Library" src="https://raw.githubusercontent.com/ARM-software/ComputeLibrary/gh-pages/ACL_logo.png" style="max-width: 100%;margin-top: 15px;margin-left: 10px"/>
   <td style="padding-left: 0.5em;">
    <div id="projectname">
    &#160;<span id="projectnumber">20.02.1</span>
    </div>
   </td>
  </tr>
  </tbody>
 </table>
 </div>
 <!-- end header part -->
 <!-- Generated by Doxygen 1.8.15 -->
 <script type="text/javascript">
 /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
 var searchBox = new SearchBox("searchBox", "search",false,'Search');
 /* @license-end */
 </script>
 <script type="text/javascript" src="menudata.js"></script>
 <script type="text/javascript" src="menu.js"></script>
 <script type="text/javascript">
 /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
 $(function() {
   initMenu('',true,false,'search.php','Search');
   $(document).ready(function() { init_search(); });
 });
 /* @license-end */</script>
 <div id="main-nav"></div>
 </div><!-- top -->
 <div id="side-nav" class="ui-resizable side-nav-resizable">
   <div id="nav-tree">
     <div id="nav-tree-contents">
       <div id="nav-sync" class="sync"></div>
     </div>
   </div>
   <div id="splitbar" style="-moz-user-select:none;"
        class="ui-resizable-handle">
   </div>
 </div>
 <script type="text/javascript">
 /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
 $(document).ready(function(){initNavTree('pooling__layer_8cl.xhtml','');});
 /* @license-end */
 </script>
 <div id="doc-content">
 <!-- window showing the filter options -->
 <div id="MSearchSelectWindow"
      onmouseover="return searchBox.OnSearchSelectShow()"
      onmouseout="return searchBox.OnSearchSelectHide()"
      onkeydown="return searchBox.OnSearchSelectKey(event)">
 </div>

 <!-- iframe showing the search results (closed by default) -->
 <div id="MSearchResultsWindow">
 <iframe src="javascript:void(0)" frameborder="0"
         name="MSearchResults" id="MSearchResults">
 </iframe>
 </div>

 <div class="header">
   <div class="summary">
 <a href="#define-members">Macros</a> &#124;
 <a href="#func-members">Functions</a>  </div>
   <div class="headertitle">
 <div class="title">pooling_layer.cl File Reference</div>  </div>
 </div><!--header-->
 <div class="contents">
 <div class="textblock"><code>#include &quot;<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
 </div>
 <p><a href="pooling__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
 <table class="memberdecls">
 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
 Macros</h2></td></tr>
 <tr class="memitem:a482ef7d59a5f474ca126e737c7f0978a"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(x,  y)&#160;&#160;&#160;((x) + (y))</td></tr>
 <tr class="separator:a482ef7d59a5f474ca126e737c7f0978a"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:a03898439d164d74f8c35bafb67262d95"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(x,  vec_size)&#160;&#160;&#160;(x)</td></tr>
 <tr class="separator:a03898439d164d74f8c35bafb67262d95"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:a5db17889d824975fefb2ce2f4690637f"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a>(x,  y)&#160;&#160;&#160;(x * (1.f / y))</td></tr>
 <tr class="separator:a5db17889d824975fefb2ce2f4690637f"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:ac9af19bec38fe50b4b9585c0e5c0ccca"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a>(x)&#160;&#160;&#160;sqrt((x))</td></tr>
 <tr class="separator:ac9af19bec38fe50b4b9585c0e5c0ccca"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:a6c01fa98d360a9d52926dc6a5a599711"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a6c01fa98d360a9d52926dc6a5a599711">DIV_OP_NHWC</a>(x,  y)&#160;&#160;&#160;(x * (<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 8))(1.f / y))</td></tr>
 <tr class="separator:a6c01fa98d360a9d52926dc6a5a599711"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:af6c2e106d0b5d1cb755bbc79d0b09d52"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(n,  <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>,  ptr)&#160;&#160;&#160;vload##n(<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)</td></tr>
 <tr class="separator:af6c2e106d0b5d1cb755bbc79d0b09d52"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:a6b8e66069b8cd3e743141d7f024a7d76"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a6b8e66069b8cd3e743141d7f024a7d76">POOLING3x3_STRIDE1</a>(res,  input,  output)</td></tr>
 <tr class="separator:a6b8e66069b8cd3e743141d7f024a7d76"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:abb8f7128361a6a1965b1b2a5b3a719b2"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#abb8f7128361a6a1965b1b2a5b3a719b2">POOLING3x3_STRIDE2</a>(res,  input,  output)</td></tr>
 <tr class="separator:abb8f7128361a6a1965b1b2a5b3a719b2"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:aba30215e4df370ff8935f83046e696ea"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#aba30215e4df370ff8935f83046e696ea">POOLING3x3_STRIDE3</a>(res,  input,  output)</td></tr>
 <tr class="separator:aba30215e4df370ff8935f83046e696ea"><td class="memSeparator" colspan="2">&#160;</td></tr>
 </table><table class="memberdecls">
 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
 Functions</h2></td></tr>
 <tr class="memitem:ae9df7602479b001f6280f2f528f10a92"><td class="memItemLeft" align="right" valign="top">ACC_DATA_TYPE&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#ae9df7602479b001f6280f2f528f10a92">calculate_avg_scale</a> (const int pool_size_x, const int pool_size_y, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y)</td></tr>
 <tr class="separator:ae9df7602479b001f6280f2f528f10a92"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:a2d95de36199fd06803ffb62f5ff1df08"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a2d95de36199fd06803ffb62f5ff1df08">pooling_layer_2</a> (__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_offset_first_element_in_bytes)</td></tr>
 <tr class="memdesc:a2d95de36199fd06803ffb62f5ff1df08"><td class="mdescLeft">&#160;</td><td class="mdescRight">Performs a pooling function of pool size equal to 2.  <a href="#a2d95de36199fd06803ffb62f5ff1df08">More...</a><br /></td></tr>
 <tr class="separator:a2d95de36199fd06803ffb62f5ff1df08"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:af5751970a4d8c62febdc6c63d6d4fd1d"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#af5751970a4d8c62febdc6c63d6d4fd1d">pooling_layer_3</a> (__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_offset_first_element_in_bytes)</td></tr>
 <tr class="memdesc:af5751970a4d8c62febdc6c63d6d4fd1d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Performs a pooling function of pool size equal to 3.  <a href="#af5751970a4d8c62febdc6c63d6d4fd1d">More...</a><br /></td></tr>
 <tr class="separator:af5751970a4d8c62febdc6c63d6d4fd1d"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:a3c82c9d144f57ed5694299ec248baad6"><td class="memItemLeft" align="right" valign="top">ACC_DATA_TYPE&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a3c82c9d144f57ed5694299ec248baad6">calculate_avg_scale_nhwc</a> (const int pool_size_x, const int pool_size_y, int upper_bound_w, int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y)</td></tr>
 <tr class="separator:a3c82c9d144f57ed5694299ec248baad6"><td class="memSeparator" colspan="2">&#160;</td></tr>
 <tr class="memitem:a074db9113f7fb9fc3f5e389892b38d32"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="pooling__layer_8cl.xhtml#a074db9113f7fb9fc3f5e389892b38d32">pooling_layer_MxN_nhwc</a> (__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_stride_w, uint input_step_w, uint input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_stride_w, uint output_step_w, uint output_offset_first_element_in_bytes)</td></tr>
 <tr class="memdesc:a074db9113f7fb9fc3f5e389892b38d32"><td class="mdescLeft">&#160;</td><td class="mdescRight">Performs a pooling function of pool size equal to N (NHWC)  <a href="#a074db9113f7fb9fc3f5e389892b38d32">More...</a><br /></td></tr>
 <tr class="separator:a074db9113f7fb9fc3f5e389892b38d32"><td class="memSeparator" colspan="2">&#160;</td></tr>
 </table>
 <h2 class="groupheader">Macro Definition Documentation</h2>
 <a id="a5db17889d824975fefb2ce2f4690637f"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a5db17889d824975fefb2ce2f4690637f">&#9670;&nbsp;</a></span>DIV_OP</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define DIV_OP</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">x, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">y&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td>&#160;&#160;&#160;(x * (1.f / y))</td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00038">38</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="a6c01fa98d360a9d52926dc6a5a599711"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a6c01fa98d360a9d52926dc6a5a599711">&#9670;&nbsp;</a></span>DIV_OP_NHWC</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define DIV_OP_NHWC</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">x, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">y&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td>&#160;&#160;&#160;(x * (<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 8))(1.f / y))</td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00041">41</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="a482ef7d59a5f474ca126e737c7f0978a"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a482ef7d59a5f474ca126e737c7f0978a">&#9670;&nbsp;</a></span>POOL_OP</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define POOL_OP</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">x, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">y&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td>&#160;&#160;&#160;((x) + (y))</td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00027">27</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="a6b8e66069b8cd3e743141d7f024a7d76"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a6b8e66069b8cd3e743141d7f024a7d76">&#9670;&nbsp;</a></span>POOLING3x3_STRIDE1</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define POOLING3x3_STRIDE1</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">res, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">input, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">output&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00059">59</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="abb8f7128361a6a1965b1b2a5b3a719b2"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#abb8f7128361a6a1965b1b2a5b3a719b2">&#9670;&nbsp;</a></span>POOLING3x3_STRIDE2</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define POOLING3x3_STRIDE2</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">res, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">input, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">output&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00102">102</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="aba30215e4df370ff8935f83046e696ea"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#aba30215e4df370ff8935f83046e696ea">&#9670;&nbsp;</a></span>POOLING3x3_STRIDE3</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define POOLING3x3_STRIDE3</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">res, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">input, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">output&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">
 <b>Value:</b><div class="fragment"><div class="line">({                                                                                                                                       \</div><div class="line">        VEC_DATA_TYPE(ACC_DATA_TYPE, 8)                                                                                                      \</div><div class="line">        data00 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(8, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 0, 0));                           \</div><div class="line">        VEC_DATA_TYPE(ACC_DATA_TYPE, 4)                                                                                                      \</div><div class="line">        data01 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(4, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 0, 0) + 8);                       \</div><div class="line">        VEC_DATA_TYPE(ACC_DATA_TYPE, 8)                                                                                                      \</div><div class="line">        data10 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(8, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 1, 0));                           \</div><div class="line">        VEC_DATA_TYPE(ACC_DATA_TYPE, 4)                                                                                                      \</div><div class="line">        data11 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(4, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 1, 0) + 8);                       \</div><div class="line">        VEC_DATA_TYPE(ACC_DATA_TYPE, 8)                                                                                                      \</div><div class="line">        data20 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(8, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 2, 0));                           \</div><div class="line">        VEC_DATA_TYPE(ACC_DATA_TYPE, 4)                                                                                                      \</div><div class="line">        data21 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(4, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 2, 0) + 8);                       \</div><div class="line">        data00 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data00, 8);                                                                                                         \</div><div class="line">        data01 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data01, 4);                                                                                                         \</div><div class="line">        data10 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data10, 8);                                                                                                         \</div><div class="line">        data11 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data11, 4);                                                                                                         \</div><div class="line">        data20 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data20, 8);                                                                                                         \</div><div class="line">        data21 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data21, 4);                                                                                                         \</div><div class="line">        \</div><div class="line">        data00 = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data00, data10);                                                                                                    \</div><div class="line">        data01 = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data01, data11);                                                                                                    \</div><div class="line">        data00 = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data00, data20);                                                                                                    \</div><div class="line">        data01 = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data01, data21);                                                                                                    \</div><div class="line">        \</div><div class="line">        res = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>((<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 4))(data00.s036, data01.s1), (<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 4))(data00.s147, data01.s2)); \</div><div class="line">        res = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(res, (<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 4))(data00.s25, data01.s03));                                                       \</div><div class="line">    })</div><div class="ttc" id="pooling__layer_8cl_xhtml_a482ef7d59a5f474ca126e737c7f0978a"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a></div><div class="ttdeci">#define POOL_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00027">pooling_layer.cl:27</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_af6c2e106d0b5d1cb755bbc79d0b09d52"><div class="ttname"><a href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a></div><div class="ttdeci">#define VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE(n, offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00056">pooling_layer.cl:56</a></div></div>
 <div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a8fcf2ddd9a1d58b1b280f5c0aed71845"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">arm_compute::test::validation::input</a></div><div class="ttdeci">auto input</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_l_s_t_m_layer_quantized_8cpp_source.xhtml#l00487">LSTMLayerQuantized.cpp:487</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_a03898439d164d74f8c35bafb67262d95"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a></div><div class="ttdeci">#define POW2_OP(x, vec_size)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00035">pooling_layer.cl:35</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00522">helpers.h:522</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">helpers.h:255</a></div></div>
 </div><!-- fragment -->
 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00142">142</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="a03898439d164d74f8c35bafb67262d95"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a03898439d164d74f8c35bafb67262d95">&#9670;&nbsp;</a></span>POW2_OP</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define POW2_OP</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">x, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">vec_size&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td>&#160;&#160;&#160;(x)</td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00035">35</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="ac9af19bec38fe50b4b9585c0e5c0ccca"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#ac9af19bec38fe50b4b9585c0e5c0ccca">&#9670;&nbsp;</a></span>SQRT_OP</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define SQRT_OP</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">x</td><td>)</td>
           <td>&#160;&#160;&#160;sqrt((x))</td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00039">39</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <a id="af6c2e106d0b5d1cb755bbc79d0b09d52"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#af6c2e106d0b5d1cb755bbc79d0b09d52">&#9670;&nbsp;</a></span>VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">#define VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</td>
           <td>(</td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">n, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname"><a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">&#160;</td>
           <td class="paramname">ptr&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td>&#160;&#160;&#160;vload##n(<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)</td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00056">56</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>

 </div>
 </div>
 <h2 class="groupheader">Function Documentation</h2>
 <a id="ae9df7602479b001f6280f2f528f10a92"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#ae9df7602479b001f6280f2f528f10a92">&#9670;&nbsp;</a></span>calculate_avg_scale()</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">ACC_DATA_TYPE calculate_avg_scale </td>
           <td>(</td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pool_size_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pool_size_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>upper_bound_w</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>upper_bound_h</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pad_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pad_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>stride_y</em>&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00172">172</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>
 <div class="fragment"><div class="line"><a name="l00174"></a><span class="lineno">  174</span>&#160;{</div><div class="line"><a name="l00175"></a><span class="lineno">  175</span>&#160;    <span class="keywordtype">int</span>       start_x = get_global_id(0) * stride_x - pad_x;</div><div class="line"><a name="l00176"></a><span class="lineno">  176</span>&#160;    <span class="keywordtype">int</span>       start_y = get_global_id(1) * stride_y - pad_y;</div><div class="line"><a name="l00177"></a><span class="lineno">  177</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> end_x   = min(start_x + pool_size_x, upper_bound_w);</div><div class="line"><a name="l00178"></a><span class="lineno">  178</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> end_y   = min(start_y + pool_size_y, upper_bound_h);</div><div class="line"><a name="l00179"></a><span class="lineno">  179</span>&#160;<span class="preprocessor">#if defined(EXCLUDE_PADDING)</span></div><div class="line"><a name="l00180"></a><span class="lineno">  180</span>&#160;    start_x = max(0, start_x);</div><div class="line"><a name="l00181"></a><span class="lineno">  181</span>&#160;    start_y = max(0, start_y);</div><div class="line"><a name="l00182"></a><span class="lineno">  182</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(EXCLUDE_PADDING) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00183"></a><span class="lineno">  183</span>&#160;    <span class="keywordflow">return</span> ((end_y - start_y) * (end_x - start_x));</div><div class="line"><a name="l00184"></a><span class="lineno">  184</span>&#160;}</div></div><!-- fragment -->
 <p class="reference">Referenced by <a class="el" href="pooling__layer_8cl_source.xhtml#l00212">pooling_layer_2()</a>, and <a class="el" href="pooling__layer_8cl_source.xhtml#l00276">pooling_layer_3()</a>.</p>

 </div>
 </div>
 <a id="a3c82c9d144f57ed5694299ec248baad6"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a3c82c9d144f57ed5694299ec248baad6">&#9670;&nbsp;</a></span>calculate_avg_scale_nhwc()</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">ACC_DATA_TYPE calculate_avg_scale_nhwc </td>
           <td>(</td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pool_size_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pool_size_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">int&#160;</td>
           <td class="paramname"><em>upper_bound_w</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">int&#160;</td>
           <td class="paramname"><em>upper_bound_h</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pad_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>pad_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">const int&#160;</td>
           <td class="paramname"><em>stride_y</em>&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00484">484</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>
 <div class="fragment"><div class="line"><a name="l00486"></a><span class="lineno">  486</span>&#160;{</div><div class="line"><a name="l00487"></a><span class="lineno">  487</span>&#160;    <span class="keywordtype">int</span> start_x = get_global_id(1) * stride_x - pad_x;</div><div class="line"><a name="l00488"></a><span class="lineno">  488</span>&#160;<span class="preprocessor">#if defined(DST_DEPTH)</span></div><div class="line"><a name="l00489"></a><span class="lineno">  489</span>&#160;    <span class="keywordtype">int</span> start_y = (get_global_id(2) % DST_DEPTH) * stride_y - pad_y;</div><div class="line"><a name="l00490"></a><span class="lineno">  490</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00491"></a><span class="lineno">  491</span>&#160;    <span class="keywordtype">int</span>       start_y    = get_global_id(2) * stride_y - pad_y;</div><div class="line"><a name="l00492"></a><span class="lineno">  492</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00493"></a><span class="lineno">  493</span>&#160;</div><div class="line"><a name="l00494"></a><span class="lineno">  494</span>&#160;<span class="preprocessor">#if !defined(EXCLUDE_PADDING)</span></div><div class="line"><a name="l00495"></a><span class="lineno">  495</span>&#160;    upper_bound_w += pad_x;</div><div class="line"><a name="l00496"></a><span class="lineno">  496</span>&#160;    upper_bound_h += pad_y;</div><div class="line"><a name="l00497"></a><span class="lineno">  497</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(EXCLUDE_PADDING) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00498"></a><span class="lineno">  498</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> end_x = min(start_x + pool_size_x, upper_bound_w);</div><div class="line"><a name="l00499"></a><span class="lineno">  499</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> end_y = min(start_y + pool_size_y, upper_bound_h);</div><div class="line"><a name="l00500"></a><span class="lineno">  500</span>&#160;<span class="preprocessor">#if defined(EXCLUDE_PADDING)</span></div><div class="line"><a name="l00501"></a><span class="lineno">  501</span>&#160;    start_x = max(0, start_x);</div><div class="line"><a name="l00502"></a><span class="lineno">  502</span>&#160;    start_y = max(0, start_y);</div><div class="line"><a name="l00503"></a><span class="lineno">  503</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(EXCLUDE_PADDING) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00504"></a><span class="lineno">  504</span>&#160;    <span class="keywordflow">return</span> ((end_y - start_y) * (end_x - start_x));</div><div class="line"><a name="l00505"></a><span class="lineno">  505</span>&#160;}</div></div><!-- fragment -->
 <p class="reference">Referenced by <a class="el" href="pooling__layer_8cl_source.xhtml#l00539">pooling_layer_MxN_nhwc()</a>.</p>

 </div>
 </div>
 <a id="a2d95de36199fd06803ffb62f5ff1df08"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a2d95de36199fd06803ffb62f5ff1df08">&#9670;&nbsp;</a></span>pooling_layer_2()</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">__kernel void pooling_layer_2 </td>
           <td>(</td>
           <td class="paramtype">__global uchar *&#160;</td>
           <td class="paramname"><em>input_ptr</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_offset_first_element_in_bytes</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">__global uchar *&#160;</td>
           <td class="paramname"><em>output_ptr</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_offset_first_element_in_bytes</em>&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p>Performs a pooling function of pool size equal to 2. </p>
 <dl class="section note"><dt>Note</dt><dd>Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16/F32; </dd>
 <dd>
 In case of average pooling the following information must be passed at compile time: -DPOOL_AVG or -DPOOL_L2 must be provided otherwise max pooling will be performed. -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad) -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension</dd></dl>
 <dl class="params"><dt>Parameters</dt><dd>
   <table class="params">
     <tr><td class="paramdir">[in]</td><td class="paramname">input_ptr</td><td>Pointer to the source image. Supported data types: F16/F32 </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_x</td><td>input_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_y</td><td>input_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_z</td><td>input_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_offset_first_element_in_bytes</td><td>The offset of the first element in the source image </td></tr>
     <tr><td class="paramdir">[out]</td><td class="paramname">output_ptr</td><td>Pointer to the destination image. Supported data types: same as <code>input_ptr</code> </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_x</td><td>Stride of the destination image in X dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_x</td><td>output_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_y</td><td>Stride of the destination image in Y dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_y</td><td>output_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_z</td><td>output_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_offset_first_element_in_bytes</td><td>The offset of the first element in the destination image </td></tr>
   </table>
   </dd>
 </dl>

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00212">212</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>
 <div class="fragment"><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;    <span class="comment">// Get pixels pointer</span></div><div class="line"><a name="l00217"></a><span class="lineno">  217</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>  = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>);</div><div class="line"><a name="l00218"></a><span class="lineno">  218</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> output = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(output);</div><div class="line"><a name="l00219"></a><span class="lineno">  219</span>&#160;</div><div class="line"><a name="l00220"></a><span class="lineno">  220</span>&#160;    <span class="comment">// Load data</span></div><div class="line"><a name="l00221"></a><span class="lineno">  221</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 2)</div><div class="line"><a name="l00222"></a><span class="lineno">  222</span>&#160;    data0 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(2, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 0, 0));</div><div class="line"><a name="l00223"></a><span class="lineno">  223</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 2)</div><div class="line"><a name="l00224"></a><span class="lineno">  224</span>&#160;    data1 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(2, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 1, 0));</div><div class="line"><a name="l00225"></a><span class="lineno">  225</span>&#160;</div><div class="line"><a name="l00226"></a><span class="lineno">  226</span>&#160;<span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00227"></a><span class="lineno">  227</span>&#160;    <span class="comment">// Raise to power of 2 for L2 Pooling</span></div><div class="line"><a name="l00228"></a><span class="lineno">  228</span>&#160;    data0 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data0, 2);</div><div class="line"><a name="l00229"></a><span class="lineno">  229</span>&#160;    data1 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data1, 2);</div><div class="line"><a name="l00230"></a><span class="lineno">  230</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></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="comment">// Perform calculations</span></div><div class="line"><a name="l00233"></a><span class="lineno">  233</span>&#160;    data0             = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data0, data1);</div><div class="line"><a name="l00234"></a><span class="lineno">  234</span>&#160;    ACC_DATA_TYPE res = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data0.s0, data0.s1);</div><div class="line"><a name="l00235"></a><span class="lineno">  235</span>&#160;</div><div class="line"><a name="l00236"></a><span class="lineno">  236</span>&#160;<span class="preprocessor">#if defined(POOL_AVG) || defined(POOL_L2)</span></div><div class="line"><a name="l00237"></a><span class="lineno">  237</span>&#160;    <span class="comment">// Divide by pool region in case of average or l2 pooling</span></div><div class="line"><a name="l00238"></a><span class="lineno">  238</span>&#160;    res = <a class="code" href="pooling__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a>(res, <a class="code" href="pooling__layer_8cl.xhtml#ae9df7602479b001f6280f2f528f10a92">calculate_avg_scale</a>(2, 2, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y));</div><div class="line"><a name="l00239"></a><span class="lineno">  239</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_AVG) || defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00240"></a><span class="lineno">  240</span>&#160;</div><div class="line"><a name="l00241"></a><span class="lineno">  241</span>&#160;<span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00242"></a><span class="lineno">  242</span>&#160;    <span class="comment">// Take square root of the result in L2 pooling</span></div><div class="line"><a name="l00243"></a><span class="lineno">  243</span>&#160;    res = <a class="code" href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a>(res);</div><div class="line"><a name="l00244"></a><span class="lineno">  244</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></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="comment">// Store result</span></div><div class="line"><a name="l00247"></a><span class="lineno">  247</span>&#160;    *(__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)output.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> = (<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>)res;</div><div class="line"><a name="l00248"></a><span class="lineno">  248</span>&#160;}</div><div class="ttc" id="pooling__layer_8cl_xhtml_a5db17889d824975fefb2ce2f4690637f"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a></div><div class="ttdeci">#define DIV_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00038">pooling_layer.cl:38</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_a482ef7d59a5f474ca126e737c7f0978a"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a></div><div class="ttdeci">#define POOL_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00027">pooling_layer.cl:27</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_af6c2e106d0b5d1cb755bbc79d0b09d52"><div class="ttname"><a href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a></div><div class="ttdeci">#define VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE(n, offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00056">pooling_layer.cl:56</a></div></div>
 <div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
 <div class="ttc" id="struct_tensor3_d_xhtml"><div class="ttname"><a href="struct_tensor3_d.xhtml">Tensor3D</a></div><div class="ttdoc">Structure to hold 3D tensor information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00358">helpers.h:358</a></div></div>
 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a8fcf2ddd9a1d58b1b280f5c0aed71845"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">arm_compute::test::validation::input</a></div><div class="ttdeci">auto input</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_l_s_t_m_layer_quantized_8cpp_source.xhtml#l00487">LSTMLayerQuantized.cpp:487</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_a03898439d164d74f8c35bafb67262d95"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a></div><div class="ttdeci">#define POW2_OP(x, vec_size)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00035">pooling_layer.cl:35</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_ae9df7602479b001f6280f2f528f10a92"><div class="ttname"><a href="pooling__layer_8cl.xhtml#ae9df7602479b001f6280f2f528f10a92">calculate_avg_scale</a></div><div class="ttdeci">ACC_DATA_TYPE calculate_avg_scale(const int pool_size_x, const int pool_size_y, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00172">pooling_layer.cl:172</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00326">helpers.h:326</a></div></div>
 <div class="ttc" id="struct_tensor3_d_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">Tensor3D::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00360">helpers.h:360</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_ac9af19bec38fe50b4b9585c0e5c0ccca"><div class="ttname"><a href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a></div><div class="ttdeci">#define SQRT_OP(x)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00039">pooling_layer.cl:39</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00522">helpers.h:522</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">helpers.h:255</a></div></div>
 </div><!-- fragment -->
 <p class="reference">References <a class="el" href="pooling__layer_8cl_source.xhtml#l00172">calculate_avg_scale()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00326">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00038">DIV_OP</a>, <a class="el" href="_c_l_2_l_s_t_m_layer_quantized_8cpp_source.xhtml#l00487">arm_compute::test::validation::input</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00027">POOL_OP</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00035">POW2_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00360">Tensor3D::ptr</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00039">SQRT_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00522">tensor3D_offset()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">VEC_DATA_TYPE</a>, and <a class="el" href="pooling__layer_8cl_source.xhtml#l00056">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>.</p>

 </div>
 </div>
 <a id="af5751970a4d8c62febdc6c63d6d4fd1d"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#af5751970a4d8c62febdc6c63d6d4fd1d">&#9670;&nbsp;</a></span>pooling_layer_3()</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">__kernel void pooling_layer_3 </td>
           <td>(</td>
           <td class="paramtype">__global uchar *&#160;</td>
           <td class="paramname"><em>input_ptr</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_offset_first_element_in_bytes</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">__global uchar *&#160;</td>
           <td class="paramname"><em>output_ptr</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_offset_first_element_in_bytes</em>&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p>Performs a pooling function of pool size equal to 3. </p>
 <dl class="section note"><dt>Note</dt><dd>Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16/F32; </dd>
 <dd>
 In case of average pooling the following information must be passed at compile time: -DPOOL_AVG or -DPOOL_L2 must be provided otherwise max pooling will be performed. -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad) -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension</dd></dl>
 <dl class="params"><dt>Parameters</dt><dd>
   <table class="params">
     <tr><td class="paramdir">[in]</td><td class="paramname">input_ptr</td><td>Pointer to the source image. Supported data types: F16/F32 </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_x</td><td>input_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_y</td><td>input_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_z</td><td>input_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_offset_first_element_in_bytes</td><td>The offset of the first element in the source image </td></tr>
     <tr><td class="paramdir">[out]</td><td class="paramname">output_ptr</td><td>Pointer to the destination image. Supported data types: same as <code>input_ptr</code> </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_x</td><td>Stride of the destination image in X dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_x</td><td>output_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_y</td><td>Stride of the destination image in Y dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_y</td><td>output_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_z</td><td>output_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_offset_first_element_in_bytes</td><td>The offset of the first element in the destination image </td></tr>
   </table>
   </dd>
 </dl>

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00276">276</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>
 <div class="fragment"><div class="line"><a name="l00279"></a><span class="lineno">  279</span>&#160;{</div><div class="line"><a name="l00280"></a><span class="lineno">  280</span>&#160;    <span class="comment">// Get pixels pointer</span></div><div class="line"><a name="l00281"></a><span class="lineno">  281</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>  = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>);</div><div class="line"><a name="l00282"></a><span class="lineno">  282</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> output = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(output);</div><div class="line"><a name="l00283"></a><span class="lineno">  283</span>&#160;</div><div class="line"><a name="l00284"></a><span class="lineno">  284</span>&#160;    <span class="comment">// Load data</span></div><div class="line"><a name="l00285"></a><span class="lineno">  285</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 3)</div><div class="line"><a name="l00286"></a><span class="lineno">  286</span>&#160;    data0 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(3, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 0, 0));</div><div class="line"><a name="l00287"></a><span class="lineno">  287</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 3)</div><div class="line"><a name="l00288"></a><span class="lineno">  288</span>&#160;    data1 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(3, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 1, 0));</div><div class="line"><a name="l00289"></a><span class="lineno">  289</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 3)</div><div class="line"><a name="l00290"></a><span class="lineno">  290</span>&#160;    data2 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(3, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, 2, 0));</div><div class="line"><a name="l00291"></a><span class="lineno">  291</span>&#160;</div><div class="line"><a name="l00292"></a><span class="lineno">  292</span>&#160;<span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00293"></a><span class="lineno">  293</span>&#160;    <span class="comment">// Raise to power of 2 for L2 Pooling</span></div><div class="line"><a name="l00294"></a><span class="lineno">  294</span>&#160;    data0 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data0, 3);</div><div class="line"><a name="l00295"></a><span class="lineno">  295</span>&#160;    data1 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data1, 3);</div><div class="line"><a name="l00296"></a><span class="lineno">  296</span>&#160;    data2 = <a class="code" href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a>(data2, 3);</div><div class="line"><a name="l00297"></a><span class="lineno">  297</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00298"></a><span class="lineno">  298</span>&#160;</div><div class="line"><a name="l00299"></a><span class="lineno">  299</span>&#160;    <span class="comment">// Perform calculations</span></div><div class="line"><a name="l00300"></a><span class="lineno">  300</span>&#160;    data0             = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data0, data1);</div><div class="line"><a name="l00301"></a><span class="lineno">  301</span>&#160;    data0             = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data0, data2);</div><div class="line"><a name="l00302"></a><span class="lineno">  302</span>&#160;    ACC_DATA_TYPE res = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(<a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(data0.s0, data0.s1), data0.s2);</div><div class="line"><a name="l00303"></a><span class="lineno">  303</span>&#160;</div><div class="line"><a name="l00304"></a><span class="lineno">  304</span>&#160;<span class="preprocessor">#if defined(POOL_AVG) || defined(POOL_L2)</span></div><div class="line"><a name="l00305"></a><span class="lineno">  305</span>&#160;    <span class="comment">// Divide by pool region in case of average pooling</span></div><div class="line"><a name="l00306"></a><span class="lineno">  306</span>&#160;    res = <a class="code" href="pooling__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a>(res, <a class="code" href="pooling__layer_8cl.xhtml#ae9df7602479b001f6280f2f528f10a92">calculate_avg_scale</a>(3, 3, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y));</div><div class="line"><a name="l00307"></a><span class="lineno">  307</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_AVG) || defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00308"></a><span class="lineno">  308</span>&#160;</div><div class="line"><a name="l00309"></a><span class="lineno">  309</span>&#160;<span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00310"></a><span class="lineno">  310</span>&#160;    <span class="comment">// Take square root of the result in L2 pooling</span></div><div class="line"><a name="l00311"></a><span class="lineno">  311</span>&#160;    res = <a class="code" href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a>(res);</div><div class="line"><a name="l00312"></a><span class="lineno">  312</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00313"></a><span class="lineno">  313</span>&#160;</div><div class="line"><a name="l00314"></a><span class="lineno">  314</span>&#160;    <span class="comment">// Store result</span></div><div class="line"><a name="l00315"></a><span class="lineno">  315</span>&#160;    *(__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)output.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> = (<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>)res;</div><div class="line"><a name="l00316"></a><span class="lineno">  316</span>&#160;}</div><div class="ttc" id="pooling__layer_8cl_xhtml_a5db17889d824975fefb2ce2f4690637f"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a></div><div class="ttdeci">#define DIV_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00038">pooling_layer.cl:38</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_a482ef7d59a5f474ca126e737c7f0978a"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a></div><div class="ttdeci">#define POOL_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00027">pooling_layer.cl:27</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_af6c2e106d0b5d1cb755bbc79d0b09d52"><div class="ttname"><a href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a></div><div class="ttdeci">#define VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE(n, offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00056">pooling_layer.cl:56</a></div></div>
 <div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
 <div class="ttc" id="struct_tensor3_d_xhtml"><div class="ttname"><a href="struct_tensor3_d.xhtml">Tensor3D</a></div><div class="ttdoc">Structure to hold 3D tensor information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00358">helpers.h:358</a></div></div>
 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a8fcf2ddd9a1d58b1b280f5c0aed71845"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">arm_compute::test::validation::input</a></div><div class="ttdeci">auto input</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_l_s_t_m_layer_quantized_8cpp_source.xhtml#l00487">LSTMLayerQuantized.cpp:487</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_a03898439d164d74f8c35bafb67262d95"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a03898439d164d74f8c35bafb67262d95">POW2_OP</a></div><div class="ttdeci">#define POW2_OP(x, vec_size)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00035">pooling_layer.cl:35</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_ae9df7602479b001f6280f2f528f10a92"><div class="ttname"><a href="pooling__layer_8cl.xhtml#ae9df7602479b001f6280f2f528f10a92">calculate_avg_scale</a></div><div class="ttdeci">ACC_DATA_TYPE calculate_avg_scale(const int pool_size_x, const int pool_size_y, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00172">pooling_layer.cl:172</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00326">helpers.h:326</a></div></div>
 <div class="ttc" id="struct_tensor3_d_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">Tensor3D::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00360">helpers.h:360</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_ac9af19bec38fe50b4b9585c0e5c0ccca"><div class="ttname"><a href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a></div><div class="ttdeci">#define SQRT_OP(x)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00039">pooling_layer.cl:39</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00522">helpers.h:522</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">helpers.h:255</a></div></div>
 </div><!-- fragment -->
 <p class="reference">References <a class="el" href="pooling__layer_8cl_source.xhtml#l00172">calculate_avg_scale()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00326">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00038">DIV_OP</a>, <a class="el" href="_c_l_2_l_s_t_m_layer_quantized_8cpp_source.xhtml#l00487">arm_compute::test::validation::input</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00027">POOL_OP</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00035">POW2_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00360">Tensor3D::ptr</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00039">SQRT_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00522">tensor3D_offset()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">VEC_DATA_TYPE</a>, and <a class="el" href="pooling__layer_8cl_source.xhtml#l00056">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>.</p>

 </div>
 </div>
 <a id="a074db9113f7fb9fc3f5e389892b38d32"></a>
 <h2 class="memtitle"><span class="permalink"><a href="#a074db9113f7fb9fc3f5e389892b38d32">&#9670;&nbsp;</a></span>pooling_layer_MxN_nhwc()</h2>

 <div class="memitem">
 <div class="memproto">
       <table class="memname">
         <tr>
           <td class="memname">__kernel void pooling_layer_MxN_nhwc </td>
           <td>(</td>
           <td class="paramtype">__global uchar *&#160;</td>
           <td class="paramname"><em>input_ptr</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_stride_w</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_step_w</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>input_offset_first_element_in_bytes</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">__global uchar *&#160;</td>
           <td class="paramname"><em>output_ptr</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_x</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_y</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_z</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_stride_w</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_step_w</em>, </td>
         </tr>
         <tr>
           <td class="paramkey"></td>
           <td></td>
           <td class="paramtype">uint&#160;</td>
           <td class="paramname"><em>output_offset_first_element_in_bytes</em>&#160;</td>
         </tr>
         <tr>
           <td></td>
           <td>)</td>
           <td></td><td></td>
         </tr>
       </table>
 </div><div class="memdoc">

 <p>Performs a pooling function of pool size equal to N (NHWC) </p>
 <dl class="section note"><dt>Note</dt><dd>Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16/F32 </dd>
 <dd>
 Pool sizes must be passed using -DPOOL_SIZE_X and -DPOOL_SIZE_Y e.g. -DPOOL_SIZE_X=13; </dd>
 <dd>
 Tensors width and height must be passed at compile time using -DMAX_WIDTH and -DMAX_HEIGHT </dd>
 <dd>
 Strides must be passed at compile time using -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions </dd>
 <dd>
 Pad values must be passed at compile time using -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension </dd>
 <dd>
 In case of average pooling the following information must be passed at compile time: -DPOOL_AVG must be provided otherwise max pooling will be performed. </dd>
 <dd>
 The initial value for the pooling operation must be passed at compile time using -DINITIAL_VALUE e.g. -DINITIAL_VALUE=0</dd></dl>
 <dl class="params"><dt>Parameters</dt><dd>
   <table class="params">
     <tr><td class="paramdir">[in]</td><td class="paramname">input_ptr</td><td>Pointer to the source image. Supported data types: F16/F32 </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_x</td><td>input_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_y</td><td>input_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_z</td><td>input_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_w</td><td>Stride of the source tensor in W dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_step_w</td><td>input_stride_w * number of elements along W processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">input_offset_first_element_in_bytes</td><td>The offset of the first element in the source image </td></tr>
     <tr><td class="paramdir">[out]</td><td class="paramname">output_ptr</td><td>Pointer to the destination image. Supported data types: same as <code>input_ptr</code> </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_x</td><td>Stride of the destination tensor in X dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_x</td><td>output_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_y</td><td>Stride of the destination tensor in Y dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_y</td><td>output_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_z</td><td>output_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_w</td><td>Stride of the destination tensor in W dimension (in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_step_w</td><td>output_stride_w * number of elements along W processed per workitem(in bytes) </td></tr>
     <tr><td class="paramdir">[in]</td><td class="paramname">output_offset_first_element_in_bytes</td><td>The offset of the first element in the destination image </td></tr>
   </table>
   </dd>
 </dl>

 <p class="definition">Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00539">539</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>
 <div class="fragment"><div class="line"><a name="l00542"></a><span class="lineno">  542</span>&#160;{</div><div class="line"><a name="l00543"></a><span class="lineno">  543</span>&#160;    <span class="comment">// Get pixels pointer</span></div><div class="line"><a name="l00544"></a><span class="lineno">  544</span>&#160;<span class="preprocessor">#if defined(DST_DEPTH)</span></div><div class="line"><a name="l00545"></a><span class="lineno">  545</span>&#160;    <a class="code" href="struct_tensor4_d.xhtml">Tensor4D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>  = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a23b9032d1b9d59547545e457f82ee478">CONVERT_TO_TENSOR4D_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, DST_DEPTH);</div><div class="line"><a name="l00546"></a><span class="lineno">  546</span>&#160;    <a class="code" href="struct_tensor4_d.xhtml">Tensor4D</a> output = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a23b9032d1b9d59547545e457f82ee478">CONVERT_TO_TENSOR4D_STRUCT</a>(output, DST_DEPTH);</div><div class="line"><a name="l00547"></a><span class="lineno">  547</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00548"></a><span class="lineno">  548</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>      = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>);</div><div class="line"><a name="l00549"></a><span class="lineno">  549</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a>  output     = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(output);</div><div class="line"><a name="l00550"></a><span class="lineno">  550</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00551"></a><span class="lineno">  551</span>&#160;</div><div class="line"><a name="l00552"></a><span class="lineno">  552</span>&#160;    <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 8)</div><div class="line"><a name="l00553"></a><span class="lineno">  553</span>&#160;    vdata = INITIAL_VALUE;</div><div class="line"><a name="l00554"></a><span class="lineno">  554</span>&#160;</div><div class="line"><a name="l00555"></a><span class="lineno">  555</span>&#160;    const <span class="keywordtype">int</span> idx_width = get_global_id(1) * STRIDE_X;</div><div class="line"><a name="l00556"></a><span class="lineno">  556</span>&#160;<span class="preprocessor">#if defined(DST_DEPTH)</span></div><div class="line"><a name="l00557"></a><span class="lineno">  557</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> idx_height = (get_global_id(2) % DST_DEPTH) * STRIDE_Y;</div><div class="line"><a name="l00558"></a><span class="lineno">  558</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00559"></a><span class="lineno">  559</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> idx_height = get_global_id(2) * STRIDE_Y;</div><div class="line"><a name="l00560"></a><span class="lineno">  560</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00561"></a><span class="lineno">  561</span>&#160;</div><div class="line"><a name="l00562"></a><span class="lineno">  562</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">int</span> y = 0; y &lt; POOL_SIZE_Y; ++y)</div><div class="line"><a name="l00563"></a><span class="lineno">  563</span>&#160;    {</div><div class="line"><a name="l00564"></a><span class="lineno">  564</span>&#160;        <span class="keywordtype">int</span> y1 = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(y, PAD_Y - idx_height, y + idx_height - PAD_Y &lt; 0 || y + idx_height - PAD_Y &gt;= MAX_HEIGHT);</div><div class="line"><a name="l00565"></a><span class="lineno">  565</span>&#160;        <span class="keywordflow">for</span>(<span class="keywordtype">int</span> x = 0; x &lt; POOL_SIZE_X; ++x)</div><div class="line"><a name="l00566"></a><span class="lineno">  566</span>&#160;        {</div><div class="line"><a name="l00567"></a><span class="lineno">  567</span>&#160;            <span class="keywordtype">int</span> x1 = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(x, PAD_X - idx_width - 1, x + idx_width - PAD_X &lt; 0 || x + idx_width - PAD_X &gt;= MAX_WIDTH);</div><div class="line"><a name="l00568"></a><span class="lineno">  568</span>&#160;            x1     = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(x1, PAD_X - idx_width - 1, y != y1);</div><div class="line"><a name="l00569"></a><span class="lineno">  569</span>&#160;</div><div class="line"><a name="l00570"></a><span class="lineno">  570</span>&#160;<span class="preprocessor">#if defined(DST_DEPTH)</span></div><div class="line"><a name="l00571"></a><span class="lineno">  571</span>&#160;            <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 8)</div><div class="line"><a name="l00572"></a><span class="lineno">  572</span>&#160;            data0 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(8, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#ad442fb5ec8be1fff97f543150de5d822">tensor4D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, x1 - PAD_X, y1 - PAD_Y, 0));</div><div class="line"><a name="l00573"></a><span class="lineno">  573</span>&#160;<span class="preprocessor">#else  </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00574"></a><span class="lineno">  574</span>&#160;            <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 8)</div><div class="line"><a name="l00575"></a><span class="lineno">  575</span>&#160;            data0 = <a class="code" href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>(8, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">input</a>, 0, x1 - PAD_X, y1 - PAD_Y));</div><div class="line"><a name="l00576"></a><span class="lineno">  576</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(DST_DEPTH) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00577"></a><span class="lineno">  577</span>&#160;</div><div class="line"><a name="l00578"></a><span class="lineno">  578</span>&#160;<span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00579"></a><span class="lineno">  579</span>&#160;            <span class="comment">// Raise to power of 2 for L2 Pooling</span></div><div class="line"><a name="l00580"></a><span class="lineno">  580</span>&#160;            data0 *= data0;</div><div class="line"><a name="l00581"></a><span class="lineno">  581</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00582"></a><span class="lineno">  582</span>&#160;            vdata = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(vdata, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(data0, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(ACC_DATA_TYPE, 8)));</div><div class="line"><a name="l00583"></a><span class="lineno">  583</span>&#160;        }</div><div class="line"><a name="l00584"></a><span class="lineno">  584</span>&#160;    }</div><div class="line"><a name="l00585"></a><span class="lineno">  585</span>&#160;</div><div class="line"><a name="l00586"></a><span class="lineno">  586</span>&#160;<span class="preprocessor">#if defined(POOL_AVG) || defined(POOL_L2)</span></div><div class="line"><a name="l00587"></a><span class="lineno">  587</span>&#160;    <span class="comment">// Divide by pool region in case of average pooling</span></div><div class="line"><a name="l00588"></a><span class="lineno">  588</span>&#160;    vdata = <a class="code" href="pooling__layer_8cl.xhtml#a6c01fa98d360a9d52926dc6a5a599711">DIV_OP_NHWC</a>(vdata, <a class="code" href="pooling__layer_8cl.xhtml#a3c82c9d144f57ed5694299ec248baad6">calculate_avg_scale_nhwc</a>(POOL_SIZE_X, POOL_SIZE_Y, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y));</div><div class="line"><a name="l00589"></a><span class="lineno">  589</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_AVG) || defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00590"></a><span class="lineno">  590</span>&#160;</div><div class="line"><a name="l00591"></a><span class="lineno">  591</span>&#160;<span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00592"></a><span class="lineno">  592</span>&#160;    <span class="comment">// Take square root of the result in L2 pooling</span></div><div class="line"><a name="l00593"></a><span class="lineno">  593</span>&#160;    vdata = <a class="code" href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a>(vdata);</div><div class="line"><a name="l00594"></a><span class="lineno">  594</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00595"></a><span class="lineno">  595</span>&#160;</div><div class="line"><a name="l00596"></a><span class="lineno">  596</span>&#160;    <span class="comment">// Store result</span></div><div class="line"><a name="l00597"></a><span class="lineno">  597</span>&#160;    vstore8(<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(vdata, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 8)), 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)output.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00598"></a><span class="lineno">  598</span>&#160;}</div><div class="ttc" id="pooling__layer_8cl_xhtml_a482ef7d59a5f474ca126e737c7f0978a"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a></div><div class="ttdeci">#define POOL_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00027">pooling_layer.cl:27</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00261">helpers.h:261</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_af6c2e106d0b5d1cb755bbc79d0b09d52"><div class="ttname"><a href="pooling__layer_8cl.xhtml#af6c2e106d0b5d1cb755bbc79d0b09d52">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a></div><div class="ttdeci">#define VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE(n, offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00056">pooling_layer.cl:56</a></div></div>
 <div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
 <div class="ttc" id="struct_tensor3_d_xhtml"><div class="ttname"><a href="struct_tensor3_d.xhtml">Tensor3D</a></div><div class="ttdoc">Structure to hold 3D tensor information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00358">helpers.h:358</a></div></div>
 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a8fcf2ddd9a1d58b1b280f5c0aed71845"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a8fcf2ddd9a1d58b1b280f5c0aed71845">arm_compute::test::validation::input</a></div><div class="ttdeci">auto input</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_l_s_t_m_layer_quantized_8cpp_source.xhtml#l00487">LSTMLayerQuantized.cpp:487</a></div></div>
 <div class="ttc" id="struct_tensor4_d_xhtml"><div class="ttname"><a href="struct_tensor4_d.xhtml">Tensor4D</a></div><div class="ttdoc">Structure to hold 4D tensor information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00368">helpers.h:368</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_ad442fb5ec8be1fff97f543150de5d822"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#ad442fb5ec8be1fff97f543150de5d822">tensor4D_offset</a></div><div class="ttdeci">__global const uchar * tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)</div><div class="ttdoc">Get the pointer position of a Tensor4D.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00535">helpers.h:535</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a23b9032d1b9d59547545e457f82ee478"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a23b9032d1b9d59547545e457f82ee478">CONVERT_TO_TENSOR4D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00333">helpers.h:333</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00326">helpers.h:326</a></div></div>
 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_af77145fbdc6b0c8931148f5597d9de53"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">arm_compute::test::validation::select</a></div><div class="ttdeci">CLSelect select</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_select_8cpp_source.xhtml#l00164">Select.cpp:164</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_a6c01fa98d360a9d52926dc6a5a599711"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a6c01fa98d360a9d52926dc6a5a599711">DIV_OP_NHWC</a></div><div class="ttdeci">#define DIV_OP_NHWC(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00041">pooling_layer.cl:41</a></div></div>
 <div class="ttc" id="struct_tensor3_d_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">Tensor3D::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00360">helpers.h:360</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_ac9af19bec38fe50b4b9585c0e5c0ccca"><div class="ttname"><a href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a></div><div class="ttdeci">#define SQRT_OP(x)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00039">pooling_layer.cl:39</a></div></div>
 <div class="ttc" id="pooling__layer_8cl_xhtml_a3c82c9d144f57ed5694299ec248baad6"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a3c82c9d144f57ed5694299ec248baad6">calculate_avg_scale_nhwc</a></div><div class="ttdeci">ACC_DATA_TYPE calculate_avg_scale_nhwc(const int pool_size_x, const int pool_size_y, int upper_bound_w, int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y)</div><div class="ttdef"><b>Definition:</b> <a href="pooling__layer_8cl_source.xhtml#l00484">pooling_layer.cl:484</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00522">helpers.h:522</a></div></div>
 <div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">helpers.h:255</a></div></div>
 </div><!-- fragment -->
 <p class="reference">References <a class="el" href="pooling__layer_8cl_source.xhtml#l00484">calculate_avg_scale_nhwc()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00261">CONVERT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00326">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00333">CONVERT_TO_TENSOR4D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00041">DIV_OP_NHWC</a>, <a class="el" href="_c_l_2_l_s_t_m_layer_quantized_8cpp_source.xhtml#l00487">arm_compute::test::validation::input</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00027">POOL_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00360">Tensor3D::ptr</a>, <a class="el" href="_c_l_2_select_8cpp_source.xhtml#l00164">arm_compute::test::validation::select</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00039">SQRT_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00522">tensor3D_offset()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00535">tensor4D_offset()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">VEC_DATA_TYPE</a>, and <a class="el" href="pooling__layer_8cl_source.xhtml#l00056">VLOAD_AND_CONVERT_TO_ACC_DATA_TYPE</a>.</p>

 </div>
 </div>
 </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_68267d1309a1af8e8297ef4c3efbcdba.xhtml">src</a></li><li class="navelem"><a class="el" href="dir_aebb8dcc11953d78e620bbef0b9e2183.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_8c278f79c760e5c5fbd911f9870614c1.xhtml">CL</a></li><li class="navelem"><a class="el" href="dir_25885286e9dad4fa105b7b25a8031bbf.xhtml">cl_kernels</a></li><li class="navelem"><a class="el" href="pooling__layer_8cl.xhtml">pooling_layer.cl</a></li>
     <li class="footer">Generated on Thu Mar 5 2020 16:07:07 for Compute Library by
     <a href="http://www.doxygen.org/index.html">
     <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.15 </li>
   </ul>
 </div>
 </body>
 </html>