arm_compute v19.11
diff --git a/documentation/pad__layer_8cl_source.xhtml b/documentation/pad__layer_8cl_source.xhtml
new file mode 100644
index 0000000..43463ee
--- /dev/null
+++ b/documentation/pad__layer_8cl_source.xhtml
@@ -0,0 +1,130 @@
+<!-- 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/pad_layer.cl Source File</title>
+<link href="tabs.css" rel="stylesheet" type="text/css"/>
+<script type="text/javascript" src="jquery.js"></script>
+<script type="text/javascript" src="dynsections.js"></script>
+<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&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">
+  <span id="projectnumber">19.11</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&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&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&dn=gpl-2.0.txt GPL-v2 */
+$(document).ready(function(){initNavTree('pad__layer_8cl_source.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="headertitle">
+<div class="title">pad_layer.cl</div> </div>
+</div><!--header-->
+<div class="contents">
+<a href="pad__layer_8cl.xhtml">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno"> 1</span> <span class="comment">/*</span></div><div class="line"><a name="l00002"></a><span class="lineno"> 2</span> <span class="comment"> * Copyright (c) 2019 ARM Limited.</span></div><div class="line"><a name="l00003"></a><span class="lineno"> 3</span> <span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno"> 4</span> <span class="comment"> * SPDX-License-Identifier: MIT</span></div><div class="line"><a name="l00005"></a><span class="lineno"> 5</span> <span class="comment"> *</span></div><div class="line"><a name="l00006"></a><span class="lineno"> 6</span> <span class="comment"> * Permission is hereby granted, free of charge, to any person obtaining a copy</span></div><div class="line"><a name="l00007"></a><span class="lineno"> 7</span> <span class="comment"> * of this software and associated documentation files (the "Software"), to</span></div><div class="line"><a name="l00008"></a><span class="lineno"> 8</span> <span class="comment"> * deal in the Software without restriction, including without limitation the</span></div><div class="line"><a name="l00009"></a><span class="lineno"> 9</span> <span class="comment"> * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or</span></div><div class="line"><a name="l00010"></a><span class="lineno"> 10</span> <span class="comment"> * sell copies of the Software, and to permit persons to whom the Software is</span></div><div class="line"><a name="l00011"></a><span class="lineno"> 11</span> <span class="comment"> * furnished to do so, subject to the following conditions:</span></div><div class="line"><a name="l00012"></a><span class="lineno"> 12</span> <span class="comment"> *</span></div><div class="line"><a name="l00013"></a><span class="lineno"> 13</span> <span class="comment"> * The above copyright notice and this permission notice shall be included in all</span></div><div class="line"><a name="l00014"></a><span class="lineno"> 14</span> <span class="comment"> * copies or substantial portions of the Software.</span></div><div class="line"><a name="l00015"></a><span class="lineno"> 15</span> <span class="comment"> *</span></div><div class="line"><a name="l00016"></a><span class="lineno"> 16</span> <span class="comment"> * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</span></div><div class="line"><a name="l00017"></a><span class="lineno"> 17</span> <span class="comment"> * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,</span></div><div class="line"><a name="l00018"></a><span class="lineno"> 18</span> <span class="comment"> * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE</span></div><div class="line"><a name="l00019"></a><span class="lineno"> 19</span> <span class="comment"> * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER</span></div><div class="line"><a name="l00020"></a><span class="lineno"> 20</span> <span class="comment"> * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,</span></div><div class="line"><a name="l00021"></a><span class="lineno"> 21</span> <span class="comment"> * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE</span></div><div class="line"><a name="l00022"></a><span class="lineno"> 22</span> <span class="comment"> * SOFTWARE.</span></div><div class="line"><a name="l00023"></a><span class="lineno"> 23</span> <span class="comment"> */</span></div><div class="line"><a name="l00024"></a><span class="lineno"> 24</span> <span class="preprocessor">#include "<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml">helpers.h</a>"</span></div><div class="line"><a name="l00025"></a><span class="lineno"> 25</span> </div><div class="line"><a name="l00026"></a><span class="lineno"> 26</span> <span class="preprocessor">#if defined(DATA_TYPE) && defined(SELECT_DT) && defined(VEC_SIZE) && defined(PAD_X_BEFORE) && defined(SRC_WIDTH)</span></div><div class="line"><a name="l00027"></a><span class="lineno"> 27</span> </div><div class="line"><a name="l00028"></a><span class="lineno"> 28</span> <span class="preprocessor">#define VEC_TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)</span></div><div class="line"><a name="l00029"></a><span class="lineno"> 29</span> <span class="preprocessor">#define VEC_INT VEC_DATA_TYPE(int, VEC_SIZE)</span></div><div class="line"><a name="l00030"></a><span class="lineno"> 30</span> <span class="preprocessor">#define VEC_SELECT VEC_DATA_TYPE(SELECT_DT, VEC_SIZE)</span></div><div class="line"><a name="l00031"></a><span class="lineno"> 31</span> <span class="preprocessor">#define OFFSETS VEC_OFFS(VEC_SELECT, VEC_SIZE)</span></div><div class="line"><a name="l00032"></a><span class="lineno"> 32</span> </div><div class="line"><a name="l00033"></a><span class="lineno"> 33</span> <span class="preprocessor">#if defined(CONST_VAL)</span></div><div class="line"><a name="l00034"></a><span class="lineno"> 34</span> <span class="comment">/** Perform a pad operation when PaddingMode is CONSTANT</span></div><div class="line"><a name="l00035"></a><span class="lineno"> 35</span> <span class="comment"> *</span></div><div class="line"><a name="l00036"></a><span class="lineno"> 36</span> <span class="comment"> * @note Data type can be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=float</span></div><div class="line"><a name="l00037"></a><span class="lineno"> 37</span> <span class="comment"> * @note Vector size must be passed using the -DVEC_SIZE compile flag, e.g. -DVEC_SIZE=4</span></div><div class="line"><a name="l00038"></a><span class="lineno"> 38</span> <span class="comment"> * @note Constant value used to fill the pads must be passed using the -DCONST_VAL compile flag, e.g. -DCONST_VAL=1.27</span></div><div class="line"><a name="l00039"></a><span class="lineno"> 39</span> <span class="comment"> * @note Pad to add to the left must be passed using the -DPAD_X_BEFORE compile flag, e.g. -DPAD_X_BEFORE=5</span></div><div class="line"><a name="l00040"></a><span class="lineno"> 40</span> <span class="comment"> * @note Input tensor's width must be passed using the -DSRC_WIDTH compile flag, e.g. -DSRC_WIDTH=224</span></div><div class="line"><a name="l00041"></a><span class="lineno"> 41</span> <span class="comment"> * @note Data type to use for the select instruction must be passed using the -DSELECT_DT compile flag, e.g. -DSELECT_DT=float</span></div><div class="line"><a name="l00042"></a><span class="lineno"> 42</span> <span class="comment"> * @note In case pad left is more than the vector size, the number of threads to skip along the X axis must be passed using the</span></div><div class="line"><a name="l00043"></a><span class="lineno"> 43</span> <span class="comment"> * -DNUM_THREADS_TO_SKIP_X compile flag, e.g. -DNUM_THREADS_TO_SKIP_X=1. This is defined as (PAD_X_BEFORE / VEC_SIZE)</span></div><div class="line"><a name="l00044"></a><span class="lineno"> 44</span> <span class="comment"> * @note If pad also needs to be added to the top of the tensor, the following compile flags must be passed at compile time:</span></div><div class="line"><a name="l00045"></a><span class="lineno"> 45</span> <span class="comment"> * -# -DPAD_Y_BEFORE: Pad to add to the top of the input tensor (e.g. -DPAD_Y_BEFORE=3)</span></div><div class="line"><a name="l00046"></a><span class="lineno"> 46</span> <span class="comment"> * -# -DSRC_HEIGHT: Input tensor's height (e.g. -DSRC_HEIGHT=127)</span></div><div class="line"><a name="l00047"></a><span class="lineno"> 47</span> <span class="comment"> * @note If pad also needs to be added to the depth of the tensor, the following compile flags must be passed at compile time:</span></div><div class="line"><a name="l00048"></a><span class="lineno"> 48</span> <span class="comment"> * -# -DPAD_Z_BEFORE: Pad to add before the first plane of the input tensor (e.g. -DPAD_Z_BEFORE=3)</span></div><div class="line"><a name="l00049"></a><span class="lineno"> 49</span> <span class="comment"> * -# -DSRC_DEPTH: Input tensor's depth (e.g. -DSRC_DEPTH=32)</span></div><div class="line"><a name="l00050"></a><span class="lineno"> 50</span> <span class="comment"> * @note If pad also needs to be added to the batch of the tensor, the following compile flags must be passed at compile time:</span></div><div class="line"><a name="l00051"></a><span class="lineno"> 51</span> <span class="comment"> * -# -DPAD_W_BEFORE: Pad to add before the first batch of the input tensor (e.g. -DPAD_W_BEFORE=3)</span></div><div class="line"><a name="l00052"></a><span class="lineno"> 52</span> <span class="comment"> * -# -DSRC_BATCH: Input tensor's batch size (e.g. -DSRC_BATCH=4)</span></div><div class="line"><a name="l00053"></a><span class="lineno"> 53</span> <span class="comment"> *</span></div><div class="line"><a name="l00054"></a><span class="lineno"> 54</span> <span class="comment"> * @param[in] src_ptr Pointer to the source image. Supported data types: U8, S8, QASYMM8, U16, S16, U32, S32, F16, F32</span></div><div class="line"><a name="l00055"></a><span class="lineno"> 55</span> <span class="comment"> * @param[in] src_stride_x Stride of the source image in X dimension (in bytes)</span></div><div class="line"><a name="l00056"></a><span class="lineno"> 56</span> <span class="comment"> * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)</span></div><div class="line"><a name="l00057"></a><span class="lineno"> 57</span> <span class="comment"> * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)</span></div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span> <span class="comment"> * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)</span></div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span> <span class="comment"> * @param[in] src_stride_z Stride of the source image in Z dimension (in bytes)</span></div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span> <span class="comment"> * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)</span></div><div class="line"><a name="l00061"></a><span class="lineno"> 61</span> <span class="comment"> * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image</span></div><div class="line"><a name="l00062"></a><span class="lineno"> 62</span> <span class="comment"> * @param[out] dst_ptr Pointer to the destination image. Supported data types: same as @p src_ptr</span></div><div class="line"><a name="l00063"></a><span class="lineno"> 63</span> <span class="comment"> * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes)</span></div><div class="line"><a name="l00064"></a><span class="lineno"> 64</span> <span class="comment"> * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)</span></div><div class="line"><a name="l00065"></a><span class="lineno"> 65</span> <span class="comment"> * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes)</span></div><div class="line"><a name="l00066"></a><span class="lineno"> 66</span> <span class="comment"> * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)</span></div><div class="line"><a name="l00067"></a><span class="lineno"> 67</span> <span class="comment"> * @param[in] dst_stride_z Stride of the destination image in Z dimension (in bytes)</span></div><div class="line"><a name="l00068"></a><span class="lineno"> 68</span> <span class="comment"> * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes)</span></div><div class="line"><a name="l00069"></a><span class="lineno"> 69</span> <span class="comment"> * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image</span></div><div class="line"><a name="l00070"></a><span class="lineno"> 70</span> <span class="comment"> * @param[in] batch (Optional) Batch index if 4D pad must be applied</span></div><div class="line"><a name="l00071"></a><span class="lineno"> 71</span> <span class="comment"> */</span></div><div class="line"><a name="l00072"></a><span class="lineno"> 72</span> __kernel <span class="keywordtype">void</span> pad_layer_constant(<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a6b83038822d1ae7ab619b684ed3b7fc0">TENSOR3D_DECLARATION</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>),</div><div class="line"><a name="l00073"></a><span class="lineno"> 73</span>  <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a6b83038822d1ae7ab619b684ed3b7fc0">TENSOR3D_DECLARATION</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>)</div><div class="line"><a name="l00074"></a><span class="lineno"> 74</span> #<span class="keywordflow">if</span> defined(PAD_W_BEFORE)</div><div class="line"><a name="l00075"></a><span class="lineno"> 75</span>  ,</div><div class="line"><a name="l00076"></a><span class="lineno"> 76</span>  uint batch</div><div class="line"><a name="l00077"></a><span class="lineno"> 77</span> #endif <span class="comment">// defined(PAD_W_BEFORE)</span></div><div class="line"><a name="l00078"></a><span class="lineno"> 78</span>  )</div><div class="line"><a name="l00079"></a><span class="lineno"> 79</span> {</div><div class="line"><a name="l00080"></a><span class="lineno"> 80</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> x = get_global_id(0);</div><div class="line"><a name="l00081"></a><span class="lineno"> 81</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> y = get_global_id(1);</div><div class="line"><a name="l00082"></a><span class="lineno"> 82</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> z = get_global_id(2);</div><div class="line"><a name="l00083"></a><span class="lineno"> 83</span> </div><div class="line"><a name="l00084"></a><span class="lineno"> 84</span>  uint cond = 0;</div><div class="line"><a name="l00085"></a><span class="lineno"> 85</span> </div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span> <span class="preprocessor">#if defined(PAD_W_BEFORE)</span></div><div class="line"><a name="l00087"></a><span class="lineno"> 87</span>  cond |= batch < PAD_W_BEFORE || batch >= (SRC_BATCH + PAD_W_BEFORE);</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span> <span class="preprocessor">#endif // defined(PAD_W_BEFORE)</span></div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span> <span class="preprocessor">#if defined(PAD_Z_BEFORE)</span></div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>  cond |= z < PAD_Z_BEFORE || z >= (SRC_DEPTH + PAD_Z_BEFORE);</div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span> <span class="preprocessor">#endif // defined(PAD_Z_BEFORE)</span></div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span> </div><div class="line"><a name="l00093"></a><span class="lineno"> 93</span>  <span class="keywordflow">if</span>(cond)</div><div class="line"><a name="l00094"></a><span class="lineno"> 94</span>  {</div><div class="line"><a name="l00095"></a><span class="lineno"> 95</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</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#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00096"></a><span class="lineno"> 96</span>  <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>)</div><div class="line"><a name="l00097"></a><span class="lineno"> 97</span>  ((VEC_TYPE)CONST_VAL, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>.ptr);</div><div class="line"><a name="l00098"></a><span class="lineno"> 98</span>  }</div><div class="line"><a name="l00099"></a><span class="lineno"> 99</span>  <span class="keywordflow">else</span></div><div class="line"><a name="l00100"></a><span class="lineno"> 100</span>  {</div><div class="line"><a name="l00101"></a><span class="lineno"> 101</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</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#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00102"></a><span class="lineno"> 102</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</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#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00103"></a><span class="lineno"> 103</span> </div><div class="line"><a name="l00104"></a><span class="lineno"> 104</span> <span class="preprocessor">#if defined(NUM_THREADS_TO_SKIP_X)</span></div><div class="line"><a name="l00105"></a><span class="lineno"> 105</span>  <span class="comment">/* In case the pad left is greater than the vector size, and we are past the threads operating solely on pad values,</span></div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span> <span class="comment"> * the input pointer must be brought back along the X axis to start from the first non-pad values.</span></div><div class="line"><a name="l00107"></a><span class="lineno"> 107</span> <span class="comment"> *</span></div><div class="line"><a name="l00108"></a><span class="lineno"> 108</span> <span class="comment"> * E.g. with VEC_SIZE=2, PAD_X_BEFORE=5, CONST_VAL=0 and 1D input |1 2 3 4 5 6|:</span></div><div class="line"><a name="l00109"></a><span class="lineno"> 109</span> <span class="comment"> * -# The first thread will compute the output values |0 0| since it detects (x_outs == (0, 1)) < PAD_X_BEFORE</span></div><div class="line"><a name="l00110"></a><span class="lineno"> 110</span> <span class="comment"> * -# The second thread will compute the output values |0 0| since it detects (x_outs == (2, 3)) < PAD_X_BEFORE</span></div><div class="line"><a name="l00111"></a><span class="lineno"> 111</span> <span class="comment"> * -# The third thread should compute |0 1|, however the input pointer is now ahead of ((x * VEC_SIZE) == 4) values, reading |4 5|</span></div><div class="line"><a name="l00112"></a><span class="lineno"> 112</span> <span class="comment"> * -# To detect this, we use ((PAD_X_BEFORE / VEC_SIZE) == NUM_THREADS_TO_SKIP_X == 2) and check that it is >= to the current x</span></div><div class="line"><a name="l00113"></a><span class="lineno"> 113</span> <span class="comment"> * -# So, we bring the pointer back of NUM_THREADS_TO_SKIP_X threads, which means multiplying this constant by the input's step along the X axis</span></div><div class="line"><a name="l00114"></a><span class="lineno"> 114</span> <span class="comment"> * -# Now that the pointer is back of ((NUM_THREADS_TO_SKIP_X * src_step_x) == 4) values, it will read the desired values |0 1|</span></div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span> <span class="comment"> */</span></div><div class="line"><a name="l00116"></a><span class="lineno"> 116</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>.ptr -= <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(0u, NUM_THREADS_TO_SKIP_X * src_step_x, x >= NUM_THREADS_TO_SKIP_X);</div><div class="line"><a name="l00117"></a><span class="lineno"> 117</span> <span class="preprocessor">#endif // defined(NUM_THREADS_TO_SKIP_X)</span></div><div class="line"><a name="l00118"></a><span class="lineno"> 118</span> <span class="preprocessor">#if defined(PAD_Z_BEFORE)</span></div><div class="line"><a name="l00119"></a><span class="lineno"> 119</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>.ptr -= PAD_Z_BEFORE * src_step_z;</div><div class="line"><a name="l00120"></a><span class="lineno"> 120</span> <span class="preprocessor">#endif // defined(PAD_Z_BEFORE)</span></div><div class="line"><a name="l00121"></a><span class="lineno"> 121</span> <span class="preprocessor">#if defined(PAD_W_BEFORE)</span></div><div class="line"><a name="l00122"></a><span class="lineno"> 122</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>.ptr -= PAD_W_BEFORE * SRC_DEPTH * src_step_z;</div><div class="line"><a name="l00123"></a><span class="lineno"> 123</span> <span class="preprocessor">#endif // defined(PAD_W_BEFORE)</span></div><div class="line"><a name="l00124"></a><span class="lineno"> 124</span> </div><div class="line"><a name="l00125"></a><span class="lineno"> 125</span>  VEC_TYPE src_vals = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>.ptr);</div><div class="line"><a name="l00126"></a><span class="lineno"> 126</span> </div><div class="line"><a name="l00127"></a><span class="lineno"> 127</span>  <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a> xs_out = (<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>)(x * <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>) + <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(OFFSETS, <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>);</div><div class="line"><a name="l00128"></a><span class="lineno"> 128</span>  <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a> cond = xs_out < (<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>)PAD_X_BEFORE || xs_out >= (<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>)(SRC_WIDTH + PAD_X_BEFORE);</div><div class="line"><a name="l00129"></a><span class="lineno"> 129</span> <span class="preprocessor">#if defined(PAD_Y_BEFORE)</span></div><div class="line"><a name="l00130"></a><span class="lineno"> 130</span>  cond |= (<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>)y < (<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>)PAD_Y_BEFORE || (<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>)y >= (<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a>)(SRC_HEIGHT + PAD_Y_BEFORE);</div><div class="line"><a name="l00131"></a><span class="lineno"> 131</span> <span class="preprocessor">#endif // defined(PAD_Y_BEFORE)</span></div><div class="line"><a name="l00132"></a><span class="lineno"> 132</span>  <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>)</div><div class="line"><a name="l00133"></a><span class="lineno"> 133</span>  (<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(src_vals, (VEC_TYPE)CONST_VAL, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(cond, VEC_SELECT)), 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>.ptr);</div><div class="line"><a name="l00134"></a><span class="lineno"> 134</span>  }</div><div class="line"><a name="l00135"></a><span class="lineno"> 135</span> }</div><div class="line"><a name="l00136"></a><span class="lineno"> 136</span> <span class="preprocessor">#endif // defined(CONST_VAL)</span></div><div class="line"><a name="l00137"></a><span class="lineno"> 137</span> </div><div class="line"><a name="l00138"></a><span class="lineno"> 138</span> <span class="preprocessor">#if defined(PAD_X_BEFORE_REMAINDER) && defined(PAD_X_AFTER_REMAINDER) && defined(PAD_X_BEFORE_REMAINDER_REFL) && defined(PAD_X_AFTER_REMAINDER_REFL) && defined(AFTER_PAD_FACT_X)</span></div><div class="line"><a name="l00139"></a><span class="lineno"> 139</span> </div><div class="line"><a name="l00140"></a><span class="lineno"> 140</span> <span class="preprocessor">#define SCALAR_COND(x) (VEC_SELECT) x == (VEC_SELECT)1</span></div><div class="line"><a name="l00141"></a><span class="lineno"> 141</span> <span class="preprocessor">#define ROTATE_REVERSE(x, n) ROTATE(REVERSE(x, VEC_SIZE), VEC_SIZE, n)</span></div><div class="line"><a name="l00142"></a><span class="lineno"> 142</span> <span class="preprocessor">#define SYMM_REFL_LEFT(x, n0, n1) select(ROTATE_REVERSE(x, n1), ROTATE(x, VEC_SIZE, n0), OFFSETS >= (VEC_SELECT)n0)</span></div><div class="line"><a name="l00143"></a><span class="lineno"> 143</span> <span class="preprocessor">#define SYMM_REFL_RIGHT(x, n0, n1) select(ROTATE(x, VEC_SIZE, n0), ROTATE_REVERSE(x, n1), OFFSETS >= (VEC_SELECT)n0)</span></div><div class="line"><a name="l00144"></a><span class="lineno"> 144</span> <span class="comment"></span></div><div class="line"><a name="l00145"></a><span class="lineno"> 145</span> <span class="comment">/** Perform a pad operation when PaddingMode is SYMMETRIC</span></div><div class="line"><a name="l00146"></a><span class="lineno"> 146</span> <span class="comment"> *</span></div><div class="line"><a name="l00147"></a><span class="lineno"> 147</span> <span class="comment"> * @note Data type can be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=float</span></div><div class="line"><a name="l00148"></a><span class="lineno"> 148</span> <span class="comment"> * @note Vector size must be passed using the -DVEC_SIZE compile flag, e.g. -DVEC_SIZE=4</span></div><div class="line"><a name="l00149"></a><span class="lineno"> 149</span> <span class="comment"> * @note Constant value must be passed using the -DCONST_VAL compile flag, e.g. -DCONST_VAL=1.27</span></div><div class="line"><a name="l00150"></a><span class="lineno"> 150</span> <span class="comment"> * @note Pad to add to the left must be passed using the -DPAD_X_BEFORE compile flag, e.g. -DPAD_X_BEFORE=5</span></div><div class="line"><a name="l00151"></a><span class="lineno"> 151</span> <span class="comment"> * @note Input tensor's width must be passed using the -DSRC_WIDTH compile flag, e.g. -DSRC_WIDTH=224</span></div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span> <span class="comment"> * @note Data type to use for the select instruction must be passed using the -DSELECT_DT compile flag, e.g. -DSELECT_DT=float</span></div><div class="line"><a name="l00153"></a><span class="lineno"> 153</span> <span class="comment"> * @note Number of values to the left when operating across left padding must be passed using the -DPAD_X_BEFORE_REMAINDER compile flag, e.g. -DPAD_X_BEFORE_REMAINDER=5</span></div><div class="line"><a name="l00154"></a><span class="lineno"> 154</span> <span class="comment"> * @note Number of values to the left when operating across right padding must be passed using the -DPAD_X_AFTER_REMAINDER compile flag, e.g. -DPAD_X_AFTER_REMAINDER=6</span></div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span> <span class="comment"> * @note To rearrange the vectors properly, (PAD_X_BEFORE_REMAINDER + 1) must be passed when mode is REFLECT using the -DPAD_X_BEFORE_REMAINDER_REFL compile flag, e.g. -DPAD_X_BEFORE_REMAINDER=6</span></div><div class="line"><a name="l00156"></a><span class="lineno"> 156</span> <span class="comment"> * @note To rearrange the vectors properly, (PAD_X_AFTER_REMAINDER - 1) must be passed using the -DPAD_X_AFTER_REMAINDER_REFL compile flag, e.g. -DPAD_X_AFTER_REMAINDER=5</span></div><div class="line"><a name="l00157"></a><span class="lineno"> 157</span> <span class="comment"> * @note When after pad X, starting point to read backward from must be passed using the -DAFTER_PAD_FACT_X compile flag, e.g. -DAFTER_PAD_FACT_X=253</span></div><div class="line"><a name="l00158"></a><span class="lineno"> 158</span> <span class="comment"> * @note If padding mode is REFLECT, the -DIS_REFLECT compile flag must be set to 1, else it must be set to 0</span></div><div class="line"><a name="l00159"></a><span class="lineno"> 159</span> <span class="comment"> * @note If pad also needs to be added to the top of the tensor, the following compile flags must be passed at compile time:</span></div><div class="line"><a name="l00160"></a><span class="lineno"> 160</span> <span class="comment"> * -# -DPAD_Y_BEFORE: Pad to add to the top of the input tensor (e.g. -DPAD_Y_BEFORE=3)</span></div><div class="line"><a name="l00161"></a><span class="lineno"> 161</span> <span class="comment"> * -# -DSRC_HEIGHT: Input tensor's height (e.g. -DSRC_HEIGHT=127)</span></div><div class="line"><a name="l00162"></a><span class="lineno"> 162</span> <span class="comment"> * @note If pad also needs to be added to the depth of the tensor, the following compile flags must be passed at compile time:</span></div><div class="line"><a name="l00163"></a><span class="lineno"> 163</span> <span class="comment"> * -# -DPAD_Z_BEFORE: Pad to add before the first plane of the input tensor (e.g. -DPAD_Z_BEFORE=3)</span></div><div class="line"><a name="l00164"></a><span class="lineno"> 164</span> <span class="comment"> * -# -DSRC_DEPTH: Input tensor's depth (e.g. -DSRC_DEPTH=32)</span></div><div class="line"><a name="l00165"></a><span class="lineno"> 165</span> <span class="comment"> * @note If the starting point to read backward from is less than the output's last element accessed in the X, the following compile flags must be passed at compile time to avoid negative offsets:</span></div><div class="line"><a name="l00166"></a><span class="lineno"> 166</span> <span class="comment"> * -# -DAFTER_PAD_REM: Defines how much to rotate the vector if the backward calculation attempted to read from a negative offset (e.g. -DAFTER_PAD_REM=3)</span></div><div class="line"><a name="l00167"></a><span class="lineno"> 167</span> <span class="comment"> *</span></div><div class="line"><a name="l00168"></a><span class="lineno"> 168</span> <span class="comment"> * @param[in] src_ptr Pointer to the source image. Supported data types: U8, S8, QASYMM8, U16, S16, U32, S32, F16, F32</span></div><div class="line"><a name="l00169"></a><span class="lineno"> 169</span> <span class="comment"> * @param[in] src_stride_x Stride of the source image in X dimension (in bytes)</span></div><div class="line"><a name="l00170"></a><span class="lineno"> 170</span> <span class="comment"> * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)</span></div><div class="line"><a name="l00171"></a><span class="lineno"> 171</span> <span class="comment"> * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)</span></div><div class="line"><a name="l00172"></a><span class="lineno"> 172</span> <span class="comment"> * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)</span></div><div class="line"><a name="l00173"></a><span class="lineno"> 173</span> <span class="comment"> * @param[in] src_stride_z Stride of the source image in Z dimension (in bytes)</span></div><div class="line"><a name="l00174"></a><span class="lineno"> 174</span> <span class="comment"> * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)</span></div><div class="line"><a name="l00175"></a><span class="lineno"> 175</span> <span class="comment"> * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image</span></div><div class="line"><a name="l00176"></a><span class="lineno"> 176</span> <span class="comment"> * @param[out] dst_ptr Pointer to the destination image. Supported data types: same as @p src_ptr</span></div><div class="line"><a name="l00177"></a><span class="lineno"> 177</span> <span class="comment"> * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes)</span></div><div class="line"><a name="l00178"></a><span class="lineno"> 178</span> <span class="comment"> * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)</span></div><div class="line"><a name="l00179"></a><span class="lineno"> 179</span> <span class="comment"> * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes)</span></div><div class="line"><a name="l00180"></a><span class="lineno"> 180</span> <span class="comment"> * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)</span></div><div class="line"><a name="l00181"></a><span class="lineno"> 181</span> <span class="comment"> * @param[in] dst_stride_z Stride of the destination image in Z dimension (in bytes)</span></div><div class="line"><a name="l00182"></a><span class="lineno"> 182</span> <span class="comment"> * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes)</span></div><div class="line"><a name="l00183"></a><span class="lineno"> 183</span> <span class="comment"> * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image</span></div><div class="line"><a name="l00184"></a><span class="lineno"> 184</span> <span class="comment"> */</span></div><div class="line"><a name="l00185"></a><span class="lineno"> 185</span> __kernel <span class="keywordtype">void</span> pad_layer_symmetric_reflect(<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a6b83038822d1ae7ab619b684ed3b7fc0">TENSOR3D_DECLARATION</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>),</div><div class="line"><a name="l00186"></a><span class="lineno"> 186</span>  <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a6b83038822d1ae7ab619b684ed3b7fc0">TENSOR3D_DECLARATION</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>))</div><div class="line"><a name="l00187"></a><span class="lineno"> 187</span> {</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>  <span class="comment">// Get current thread position</span></div><div class="line"><a name="l00189"></a><span class="lineno"> 189</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> x = get_global_id(0);</div><div class="line"><a name="l00190"></a><span class="lineno"> 190</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> y = get_global_id(1);</div><div class="line"><a name="l00191"></a><span class="lineno"> 191</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> z = get_global_id(2);</div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span> </div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span>  <span class="comment">// Define conditions based on the thread X position w.r.t. pad left and right</span></div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> x_out_first = x * <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>;</div><div class="line"><a name="l00195"></a><span class="lineno"> 195</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> x_out_last = x_out_first + <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>;</div><div class="line"><a name="l00196"></a><span class="lineno"> 196</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> is_before_pad_left = (x_out_last <= PAD_X_BEFORE);</div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> is_across_pad_left = (x_out_first < PAD_X_BEFORE) && (x_out_last > PAD_X_BEFORE);</div><div class="line"><a name="l00198"></a><span class="lineno"> 198</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> is_inside_input = (x_out_first >= PAD_X_BEFORE) && (x_out_last <= (SRC_WIDTH + PAD_X_BEFORE));</div><div class="line"><a name="l00199"></a><span class="lineno"> 199</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> is_across_pad_right = (x_out_first < (SRC_WIDTH + PAD_X_BEFORE)) && (x_out_last > (SRC_WIDTH + PAD_X_BEFORE));</div><div class="line"><a name="l00200"></a><span class="lineno"> 200</span>  <span class="keyword">const</span> <span class="keywordtype">int</span> is_after_pad_right = (x_out_first >= (SRC_WIDTH + PAD_X_BEFORE));</div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span> </div><div class="line"><a name="l00202"></a><span class="lineno"> 202</span>  <span class="comment">// Calculate base pointers</span></div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>  __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes;</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</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#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00205"></a><span class="lineno"> 205</span> </div><div class="line"><a name="l00206"></a><span class="lineno"> 206</span>  <span class="comment">// Calculate input tensor's offset based on the defined conditions</span></div><div class="line"><a name="l00207"></a><span class="lineno"> 207</span>  <span class="keywordtype">int</span> x_offset = 0;</div><div class="line"><a name="l00208"></a><span class="lineno"> 208</span>  x_offset = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(x_offset, PAD_X_BEFORE - x_out_last + IS_REFLECT, is_before_pad_left);</div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span>  x_offset = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(x_offset, x_out_first - PAD_X_BEFORE, is_inside_input);</div><div class="line"><a name="l00210"></a><span class="lineno"> 210</span>  x_offset = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(x_offset, SRC_WIDTH - <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>, is_across_pad_right);</div><div class="line"><a name="l00211"></a><span class="lineno"> 211</span>  x_offset = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(x_offset, AFTER_PAD_FACT_X - x_out_last, is_after_pad_right);</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span> </div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span> <span class="preprocessor">#if defined(AFTER_PAD_REM)</span></div><div class="line"><a name="l00214"></a><span class="lineno"> 214</span>  <span class="keywordtype">int</span> neg_offs = x_offset < 0;</div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>  x_offset = max(x_offset, 0);</div><div class="line"><a name="l00216"></a><span class="lineno"> 216</span> <span class="preprocessor">#endif // defined(AFTER_PAD_REM)</span></div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span> </div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>  <span class="comment">// Load input values from the computed offset</span></div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span>  <span class="keywordtype">int</span> y_in = y;</div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span>  <span class="keywordtype">int</span> z_in = z;</div><div class="line"><a name="l00221"></a><span class="lineno"> 221</span> <span class="preprocessor">#if defined(PAD_Y_BEFORE)</span></div><div class="line"><a name="l00222"></a><span class="lineno"> 222</span>  y_in = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(y - PAD_Y_BEFORE, PAD_Y_BEFORE - y + IS_REFLECT - 1, y < PAD_Y_BEFORE);</div><div class="line"><a name="l00223"></a><span class="lineno"> 223</span>  y_in = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(y_in, 2 * SRC_HEIGHT + PAD_Y_BEFORE - y - IS_REFLECT - 1, y >= (SRC_HEIGHT + PAD_Y_BEFORE));</div><div class="line"><a name="l00224"></a><span class="lineno"> 224</span> <span class="preprocessor">#endif // defined(PAD_Y_BEFORE)</span></div><div class="line"><a name="l00225"></a><span class="lineno"> 225</span> <span class="preprocessor">#if defined(PAD_Z_BEFORE)</span></div><div class="line"><a name="l00226"></a><span class="lineno"> 226</span>  z_in = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(z - PAD_Z_BEFORE, PAD_Z_BEFORE - z + IS_REFLECT - 1, z < PAD_Z_BEFORE);</div><div class="line"><a name="l00227"></a><span class="lineno"> 227</span>  z_in = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(z_in, 2 * SRC_DEPTH + PAD_Z_BEFORE - z - IS_REFLECT - 1, z >= (SRC_DEPTH + PAD_Z_BEFORE));</div><div class="line"><a name="l00228"></a><span class="lineno"> 228</span> <span class="preprocessor">#endif // defined(PAD_Y_BEFORE)</span></div><div class="line"><a name="l00229"></a><span class="lineno"> 229</span> </div><div class="line"><a name="l00230"></a><span class="lineno"> 230</span>  src_addr += x_offset * src_stride_x + y_in * src_step_y + z_in * src_step_z;</div><div class="line"><a name="l00231"></a><span class="lineno"> 231</span> </div><div class="line"><a name="l00232"></a><span class="lineno"> 232</span> <span class="preprocessor">#if SRC_WIDTH == 1</span></div><div class="line"><a name="l00233"></a><span class="lineno"> 233</span>  <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>)</div><div class="line"><a name="l00234"></a><span class="lineno"> 234</span>  ((VEC_TYPE)(*(__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)src_addr), 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>.ptr);</div><div class="line"><a name="l00235"></a><span class="lineno"> 235</span> <span class="preprocessor">#else // SRC_WIDTH == 1</span></div><div class="line"><a name="l00236"></a><span class="lineno"> 236</span> </div><div class="line"><a name="l00237"></a><span class="lineno"> 237</span>  VEC_TYPE src_vals = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)src_addr);</div><div class="line"><a name="l00238"></a><span class="lineno"> 238</span> </div><div class="line"><a name="l00239"></a><span class="lineno"> 239</span>  <span class="comment">// Choose rearrangement policy based on the defined conditions</span></div><div class="line"><a name="l00240"></a><span class="lineno"> 240</span>  src_vals = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(src_vals, SYMM_REFL_LEFT(src_vals, PAD_X_BEFORE_REMAINDER, PAD_X_BEFORE_REMAINDER_REFL), SCALAR_COND(is_across_pad_left));</div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>  src_vals = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(src_vals, SYMM_REFL_RIGHT(src_vals, PAD_X_AFTER_REMAINDER, PAD_X_AFTER_REMAINDER_REFL), SCALAR_COND(is_across_pad_right));</div><div class="line"><a name="l00242"></a><span class="lineno"> 242</span>  src_vals = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(src_vals, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#abcadce1c5f0c9f446d1a72cac2421121">REVERSE</a>(src_vals, <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>), SCALAR_COND((is_before_pad_left || is_after_pad_right)));</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span> <span class="preprocessor">#if defined(AFTER_PAD_REM)</span></div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span>  src_vals = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(src_vals, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa6a1974248e56bbdfb2845d985093ce7">ROTATE</a>(src_vals, <a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>, AFTER_PAD_REM), SCALAR_COND(neg_offs));</div><div class="line"><a name="l00245"></a><span class="lineno"> 245</span> <span class="preprocessor">#endif // defined(AFTER_PAD_REM)</span></div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span> </div><div class="line"><a name="l00247"></a><span class="lineno"> 247</span>  <span class="comment">// Store</span></div><div class="line"><a name="l00248"></a><span class="lineno"> 248</span>  <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a>)</div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>  (src_vals, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>.ptr);</div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span> <span class="preprocessor">#endif // SRC_WIDTH == 1</span></div><div class="line"><a name="l00251"></a><span class="lineno"> 251</span> }</div><div class="line"><a name="l00252"></a><span class="lineno"> 252</span> <span class="preprocessor">#endif // defined(PAD_X_BEFORE_REMAINDER) && defined(PAD_X_AFTER_REMAINDER) && defined(PAD_X_BEFORE_REMAINDER_REFL) && defined(PAD_X_AFTER_REMAINDER_REFL) && defined(AFTER_PAD_FACT_X)</span></div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span> <span class="preprocessor">#endif // defined(DATA_TYPE) && defined(SELECT_DT) && defined(VEC_SIZE) && defined(PAD_X_BEFORE) && defined(SRC_WIDTH)</span></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="depthwise__convolution__quantized_8cl_xhtml_a3fffea119c04c7680f2e9cf3fadf63b4"><div class="ttname"><a href="depthwise__convolution__quantized_8cl.xhtml#a3fffea119c04c7680f2e9cf3fadf63b4">VEC_SIZE</a></div><div class="ttdeci">#define VEC_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution__quantized_8cl_source.xhtml#l00031">depthwise_convolution_quantized.cl:31</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aa6a1974248e56bbdfb2845d985093ce7"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa6a1974248e56bbdfb2845d985093ce7">ROTATE</a></div><div class="ttdeci">#define ROTATE(x, s, n)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00162">helpers.h:162</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="depthwise__convolution__quantized_8cl_xhtml_aee190caf3b3571e939ac129e12c368cd"><div class="ttname"><a href="depthwise__convolution__quantized_8cl.xhtml#aee190caf3b3571e939ac129e12c368cd">VEC_INT</a></div><div class="ttdeci">#define VEC_INT</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution__quantized_8cl_source.xhtml#l00042">depthwise_convolution_quantized.cl:42</a></div></div>
+<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_absolute_difference_8cpp_source.xhtml#l00102">AbsoluteDifference.cpp:102</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_abcadce1c5f0c9f446d1a72cac2421121"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#abcadce1c5f0c9f446d1a72cac2421121">REVERSE</a></div><div class="ttdeci">#define REVERSE(x, s)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00099">helpers.h:99</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml">helpers.h</a></div></div>
+<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_acb282042d1edeeaa3cc979a206f78b54"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a></div><div class="ttdeci">#define VSTORE(size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00198">helpers.h:198</a></div></div>
+<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a287e2fc366c312b468382c95bb90f91f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a></div><div class="ttdeci">#define VLOAD(size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00195">helpers.h:195</a></div></div>
+<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a6b83038822d1ae7ab619b684ed3b7fc0"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a6b83038822d1ae7ab619b684ed3b7fc0">TENSOR3D_DECLARATION</a></div><div class="ttdeci">#define TENSOR3D_DECLARATION(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00283">helpers.h:283</a></div></div>
+<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure & src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
+</div><!-- fragment --></div><!-- contents -->
+</div><!-- doc-content -->
+<!-- start footer part -->
+<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
+ <ul>
+ <li class="navelem"><a class="el" href="dir_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="pad__layer_8cl.xhtml">pad_layer.cl</a></li>
+ <li class="footer">Generated on Thu Nov 28 2019 16:52:58 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>