+<div class="ttc" id="pooling__layer_8cl_xhtml_ac8f27d6ce33043a58fc6bd17b41f8153"><div class="ttname"><a href="pooling__layer_8cl.xhtml#ac8f27d6ce33043a58fc6bd17b41f8153">calculate_avg_scale</a></div><div class="ttdeci">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#l00188">pooling_layer.cl:188</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00315">helpers.h:315</a></div></div>
+<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
+<div class="ttc" id="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="helpers_8h_source.xhtml#l00153">helpers.h:153</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#l00061">pooling_layer.cl:61</a></div></div>
+</div><!-- fragment -->
+</div>
+</div>
+<a class="anchor" id="ac320fea94a0f786ca35fe159d57a6e49"></a>
+<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 * </td>
+ <td class="paramname"><em>input_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>input_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>input_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>input_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>input_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>input_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>input_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </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 * </td>
+ <td class="paramname"><em>output_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>output_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>output_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>output_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>output_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>output_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>output_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>output_offset_first_element_in_bytes</em> </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>
+-DFP16 must be passed at compile time if half float data type is used </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></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>Definition at line <a class="el" href="pooling__layer_8cl_source.xhtml#l00565">565</a> of file <a class="el" href="pooling__layer_8cl_source.xhtml">pooling_layer.cl</a>.</p>
+
+<p>References <a class="el" href="pooling__layer_8cl_source.xhtml#l00518">calculate_avg_scale_nhwc()</a>, <a class="el" href="helpers_8h_source.xhtml#l00119">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00065">DIV_OP_NHWC</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00049">POOL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00153">Tensor3D::ptr</a>, <a class="el" href="pooling__layer_8cl_source.xhtml#l00061">SQRT_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00315">tensor3D_offset()</a>, and <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>.</p>
+<div class="fragment"><div class="line"><a name="l00568"></a><span class="lineno"> 568</span> {</div><div class="line"><a name="l00569"></a><span class="lineno"> 569</span>  <span class="comment">// Get pixels pointer</span></div><div class="line"><a name="l00570"></a><span class="lineno"> 570</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> input = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(input);</div><div class="line"><a name="l00571"></a><span class="lineno"> 571</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> output = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(output);</div><div class="line"><a name="l00572"></a><span class="lineno"> 572</span> </div><div class="line"><a name="l00573"></a><span class="lineno"> 573</span>  <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 8)</div><div class="line"><a name="l00574"></a><span class="lineno"> 574</span>  vdata = INITIAL_VALUE;</div><div class="line"><a name="l00575"></a><span class="lineno"> 575</span>  <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> sdata = INITIAL_VALUE;</div><div class="line"><a name="l00576"></a><span class="lineno"> 576</span> </div><div class="line"><a name="l00577"></a><span class="lineno"> 577</span>  const <span class="keywordtype">int</span> idx_width = get_global_id(1) * STRIDE_X;</div><div class="line"><a name="l00578"></a><span class="lineno"> 578</span>  const <span class="keywordtype">int</span> idx_height = get_global_id(2) * STRIDE_Y;</div><div class="line"><a name="l00579"></a><span class="lineno"> 579</span> </div><div class="line"><a name="l00580"></a><span class="lineno"> 580</span>  for(<span class="keywordtype">int</span> y = 0; y < POOL_SIZE_Y; ++y)</div><div class="line"><a name="l00581"></a><span class="lineno"> 581</span>  {</div><div class="line"><a name="l00582"></a><span class="lineno"> 582</span>  <span class="keywordtype">int</span> y1 = select(y, PAD_Y - idx_height, y + idx_height < PAD_Y || y + idx_height > MAX_HEIGHT);</div><div class="line"><a name="l00583"></a><span class="lineno"> 583</span>  <span class="keywordflow">for</span>(<span class="keywordtype">int</span> x = 0; x < POOL_SIZE_X; ++x)</div><div class="line"><a name="l00584"></a><span class="lineno"> 584</span>  {</div><div class="line"><a name="l00585"></a><span class="lineno"> 585</span>  <span class="keywordtype">int</span> x1 = select(x, PAD_X - idx_width - 1, x + idx_width < PAD_X || x + idx_width > MAX_WIDTH);</div><div class="line"><a name="l00586"></a><span class="lineno"> 586</span>  x1 = select(x1, PAD_X - idx_width - 1, y != y1);</div><div class="line"><a name="l00587"></a><span class="lineno"> 587</span> </div><div class="line"><a name="l00588"></a><span class="lineno"> 588</span>  <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 8)</div><div class="line"><a name="l00589"></a><span class="lineno"> 589</span>  data0 = vload8(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&input, 0, x1 - PAD_X, y1 - PAD_Y));</div><div class="line"><a name="l00590"></a><span class="lineno"> 590</span> <span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00591"></a><span class="lineno"> 591</span>  <span class="comment">// Raise to power of 2 for L2 Pooling</span></div><div class="line"><a name="l00592"></a><span class="lineno"> 592</span>  data0 *= data0;</div><div class="line"><a name="l00593"></a><span class="lineno"> 593</span> <span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00594"></a><span class="lineno"> 594</span>  vdata = <a class="code" href="pooling__layer_8cl.xhtml#a482ef7d59a5f474ca126e737c7f0978a">POOL_OP</a>(vdata, data0);</div><div class="line"><a name="l00595"></a><span class="lineno"> 595</span>  }</div><div class="line"><a name="l00596"></a><span class="lineno"> 596</span>  }</div><div class="line"><a name="l00597"></a><span class="lineno"> 597</span> </div><div class="line"><a name="l00598"></a><span class="lineno"> 598</span> <span class="preprocessor">#if defined(POOL_AVG) || defined(POOL_L2)</span></div><div class="line"><a name="l00599"></a><span class="lineno"> 599</span>  <span class="comment">// Divide by pool region in case of average pooling</span></div><div class="line"><a name="l00600"></a><span class="lineno"> 600</span>  vdata = <a class="code" href="pooling__layer_8cl.xhtml#a6c01fa98d360a9d52926dc6a5a599711">DIV_OP_NHWC</a>(vdata, <a class="code" href="pooling__layer_8cl.xhtml#a93173b51e52924504320e874affcf881">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="l00601"></a><span class="lineno"> 601</span> <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="l00602"></a><span class="lineno"> 602</span> </div><div class="line"><a name="l00603"></a><span class="lineno"> 603</span> <span class="preprocessor">#if defined(POOL_L2)</span></div><div class="line"><a name="l00604"></a><span class="lineno"> 604</span>  <span class="comment">// Take square root of the result in L2 pooling</span></div><div class="line"><a name="l00605"></a><span class="lineno"> 605</span>  vdata = <a class="code" href="pooling__layer_8cl.xhtml#ac9af19bec38fe50b4b9585c0e5c0ccca">SQRT_OP</a>(vdata);</div><div class="line"><a name="l00606"></a><span class="lineno"> 606</span> <span class="preprocessor">#endif </span><span class="comment">/* defined(POOL_L2) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00607"></a><span class="lineno"> 607</span> </div><div class="line"><a name="l00608"></a><span class="lineno"> 608</span>  <span class="comment">// Store result</span></div><div class="line"><a name="l00609"></a><span class="lineno"> 609</span>  vstore8(vdata, 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="l00610"></a><span class="lineno"> 610</span> }</div><div class="ttc" id="pooling__layer_8cl_xhtml_a93173b51e52924504320e874affcf881"><div class="ttname"><a href="pooling__layer_8cl.xhtml#a93173b51e52924504320e874affcf881">calculate_avg_scale_nhwc</a></div><div class="ttdeci">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#l00518">pooling_layer.cl:518</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#l00049">pooling_layer.cl:49</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00119">helpers.h:119</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="helpers_8h_source.xhtml#l00151">helpers.h:151</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="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="helpers_8h_source.xhtml#l00315">helpers.h:315</a></div></div>