documentation/depthwise__convolution_8cl.xhtml

   1 <!-- HTML header for doxygen 1.8.9.1-->
   2 <!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
   3 <html xmlns="http://www.w3.org/1999/xhtml">
   4 <head>
   5 <meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
   6 <meta http-equiv="X-UA-Compatible" content="IE=9"/>
   7 <meta name="generator" content="Doxygen 1.8.11"/>
   8 <meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
   9 <title>Compute Library: src/core/CL/cl_kernels/depthwise_convolution.cl File Reference</title>
  10 <link href="tabs.css" rel="stylesheet" type="text/css"/>
  11 <script type="text/javascript" src="jquery.js"></script>
  12 <script type="text/javascript" src="dynsections.js"></script>
  13 <link href="navtree.css" rel="stylesheet" type="text/css"/>
  14 <script type="text/javascript" src="resize.js"></script>
  15 <script type="text/javascript" src="navtreedata.js"></script>
  16 <script type="text/javascript" src="navtree.js"></script>
  17 <script type="text/javascript">
  18   $(document).ready(initResizable);
  19   $(window).load(resizeHeight);
  20 </script>
  21 <link href="search/search.css" rel="stylesheet" type="text/css"/>
  22 <script type="text/javascript" src="search/searchdata.js"></script>
  23 <script type="text/javascript" src="search/search.js"></script>
  24 <script type="text/javascript">
  25   $(document).ready(function() { init_search(); });
  26 </script>
  27 <script type="text/x-mathjax-config">
  28   MathJax.Hub.Config({
  29     extensions: ["tex2jax.js"],
  30     jax: ["input/TeX","output/HTML-CSS"],
  31 });
  32 </script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
  33 <link href="doxygen.css" rel="stylesheet" type="text/css" />
  34 </head>
  35 <body>
  36 <div id="top"><!-- do not remove this div, it is closed by doxygen! -->
  37 <div id="titlearea">
  38 <table cellspacing="0" cellpadding="0">
  39  <tbody>
  40  <tr style="height: 56px;">
  41   <td style="padding-left: 0.5em;">
  42    <div id="projectname">Compute Library
  43    &#160;<span id="projectnumber">18.02</span>
  44    </div>
  45   </td>
  46  </tr>
  47  </tbody>
  48 </table>
  49 </div>
  50 <!-- end header part -->
  51 <!-- Generated by Doxygen 1.8.11 -->
  52 <script type="text/javascript">
  53 var searchBox = new SearchBox("searchBox", "search",false,'Search');
  54 </script>
  55   <div id="navrow1" class="tabs">
  56     <ul class="tablist">
  57       <li><a href="index.xhtml"><span>Main&#160;Page</span></a></li>
  58       <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
  59       <li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
  60       <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
  61       <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
  62       <li>
  63         <div id="MSearchBox" class="MSearchBoxInactive">
  64         <span class="left">
  65           <img id="MSearchSelect" src="search/mag_sel.png"
  66                onmouseover="return searchBox.OnSearchSelectShow()"
  67                onmouseout="return searchBox.OnSearchSelectHide()"
  68                alt=""/>
  69           <input type="text" id="MSearchField" value="Search" accesskey="S"
  70                onfocus="searchBox.OnSearchFieldFocus(true)"
  71                onblur="searchBox.OnSearchFieldFocus(false)"
  72                onkeyup="searchBox.OnSearchFieldChange(event)"/>
  73           </span><span class="right">
  74             <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
  75           </span>
  76         </div>
  77       </li>
  78     </ul>
  79   </div>
  80   <div id="navrow2" class="tabs2">
  81     <ul class="tablist">
  82       <li><a href="files.xhtml"><span>File&#160;List</span></a></li>
  83       <li><a href="globals.xhtml"><span>Globals</span></a></li>
  84     </ul>
  85   </div>
  86 </div><!-- top -->
  87 <div id="side-nav" class="ui-resizable side-nav-resizable">
  88   <div id="nav-tree">
  89     <div id="nav-tree-contents">
  90       <div id="nav-sync" class="sync"></div>
  91     </div>
  92   </div>
  93   <div id="splitbar" style="-moz-user-select:none;"
  94        class="ui-resizable-handle">
  95   </div>
  96 </div>
  97 <script type="text/javascript">
  98 $(document).ready(function(){initNavTree('depthwise__convolution_8cl.xhtml','');});
  99 </script>
 100 <div id="doc-content">
 101 <!-- window showing the filter options -->
 102 <div id="MSearchSelectWindow"
 103      onmouseover="return searchBox.OnSearchSelectShow()"
 104      onmouseout="return searchBox.OnSearchSelectHide()"
 105      onkeydown="return searchBox.OnSearchSelectKey(event)">
 106 </div>
 107
 108 <!-- iframe showing the search results (closed by default) -->
 109 <div id="MSearchResultsWindow">
 110 <iframe src="javascript:void(0)" frameborder="0"
 111         name="MSearchResults" id="MSearchResults">
 112 </iframe>
 113 </div>
 114
 115 <div class="header">
 116   <div class="summary">
 117 <a href="#define-members">Macros</a> &#124;
 118 <a href="#func-members">Functions</a>  </div>
 119   <div class="headertitle">
 120 <div class="title">depthwise_convolution.cl File Reference</div>  </div>
 121 </div><!--header-->
 122 <div class="contents">
 123 <div class="textblock"><code>#include &quot;<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
 124 </div>
 125 <p><a href="depthwise__convolution_8cl_source.xhtml">Go to the source code of this file.</a></p>
 126 <table class="memberdecls">
 127 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
 128 Macros</h2></td></tr>
 129 <tr class="memitem:aa5ab4cf85fe56de3821727cfcef327d9"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(acc,  src0,  weights_row0)</td></tr>
 130 <tr class="separator:aa5ab4cf85fe56de3821727cfcef327d9"><td class="memSeparator" colspan="2">&#160;</td></tr>
 131 <tr class="memitem:a31cbdd6b28ee6331839432f69d87fdcb"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(acc,  src0,  src1,  weights_row0)</td></tr>
 132 <tr class="separator:a31cbdd6b28ee6331839432f69d87fdcb"><td class="memSeparator" colspan="2">&#160;</td></tr>
 133 </table><table class="memberdecls">
 134 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
 135 Functions</h2></td></tr>
 136 <tr class="memitem:a2fbbf5b73260b66b310d4d857db5289e"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a2fbbf5b73260b66b310d4d857db5289e">depthwise_convolution_3x3_stridex1_stridey1_bifrost</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_stride_z, uint src_step_z, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes, __global uchar *weights_ptr, uint weights_stride_x, uint weights_step_x, uint weights_stride_y, uint weights_step_y, uint weights_stride_z, uint weights_step_z, uint weights_offset_first_element_in_bytes, __global uchar *biases_ptr, uint biases_stride_x, uint biases_step_x, uint biases_offset_first_element_in_bytes)</td></tr>
 137 <tr class="memdesc:a2fbbf5b73260b66b310d4d857db5289e"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 1.  <a href="#a2fbbf5b73260b66b310d4d857db5289e">More...</a><br /></td></tr>
 138 <tr class="separator:a2fbbf5b73260b66b310d4d857db5289e"><td class="memSeparator" colspan="2">&#160;</td></tr>
 139 <tr class="memitem:a328ab548c536089e8238fb33de617bec"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a328ab548c536089e8238fb33de617bec">depthwise_convolution_3x3_stridex2_stridey2_bifrost</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_stride_z, uint src_step_z, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes, __global uchar *weights_ptr, uint weights_stride_x, uint weights_step_x, uint weights_stride_y, uint weights_step_y, uint weights_stride_z, uint weights_step_z, uint weights_offset_first_element_in_bytes, __global uchar *biases_ptr, uint biases_stride_x, uint biases_step_x, uint biases_offset_first_element_in_bytes)</td></tr>
 140 <tr class="memdesc:a328ab548c536089e8238fb33de617bec"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 2.  <a href="#a328ab548c536089e8238fb33de617bec">More...</a><br /></td></tr>
 141 <tr class="separator:a328ab548c536089e8238fb33de617bec"><td class="memSeparator" colspan="2">&#160;</td></tr>
 142 </table>
 143 <h2 class="groupheader">Macro Definition Documentation</h2>
 144 <a class="anchor" id="aa5ab4cf85fe56de3821727cfcef327d9"></a>
 145 <div class="memitem">
 146 <div class="memproto">
 147       <table class="memname">
 148         <tr>
 149           <td class="memname">#define CONVOLUTION1x3_BIFROST2X1_STRIDE1</td>
 150           <td>(</td>
 151           <td class="paramtype">&#160;</td>
 152           <td class="paramname">acc, </td>
 153         </tr>
 154         <tr>
 155           <td class="paramkey"></td>
 156           <td></td>
 157           <td class="paramtype">&#160;</td>
 158           <td class="paramname">src0, </td>
 159         </tr>
 160         <tr>
 161           <td class="paramkey"></td>
 162           <td></td>
 163           <td class="paramtype">&#160;</td>
 164           <td class="paramname">weights_row0&#160;</td>
 165         </tr>
 166         <tr>
 167           <td></td>
 168           <td>)</td>
 169           <td></td><td></td>
 170         </tr>
 171       </table>
 172 </div><div class="memdoc">
 173 <b>Value:</b><div class="fragment"><div class="line">({                                                             \</div><div class="line">        acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0);            \</div><div class="line">        acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0);            \</div><div class="line">        acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0);            \</div><div class="line">        acc.s1 = fma(src0.s1, weights_row0.s0, acc.s1);            \</div><div class="line">        acc.s1 = fma(src0.s2, weights_row0.s1, acc.s1);            \</div><div class="line">        acc.s1 = fma(src0.s3, weights_row0.s2, acc.s1);            \</div><div class="line">    })</div></div><!-- fragment -->
 174 <p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00211">211</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
 175
 176 <p>Referenced by <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00263">depthwise_convolution_3x3_stridex1_stridey1_bifrost()</a>.</p>
 177
 178 </div>
 179 </div>
 180 <a class="anchor" id="a31cbdd6b28ee6331839432f69d87fdcb"></a>
 181 <div class="memitem">
 182 <div class="memproto">
 183       <table class="memname">
 184         <tr>
 185           <td class="memname">#define CONVOLUTION1x3_BIFROST2X1_STRIDE2</td>
 186           <td>(</td>
 187           <td class="paramtype">&#160;</td>
 188           <td class="paramname">acc, </td>
 189         </tr>
 190         <tr>
 191           <td class="paramkey"></td>
 192           <td></td>
 193           <td class="paramtype">&#160;</td>
 194           <td class="paramname">src0, </td>
 195         </tr>
 196         <tr>
 197           <td class="paramkey"></td>
 198           <td></td>
 199           <td class="paramtype">&#160;</td>
 200           <td class="paramname">src1, </td>
 201         </tr>
 202         <tr>
 203           <td class="paramkey"></td>
 204           <td></td>
 205           <td class="paramtype">&#160;</td>
 206           <td class="paramname">weights_row0&#160;</td>
 207         </tr>
 208         <tr>
 209           <td></td>
 210           <td>)</td>
 211           <td></td><td></td>
 212         </tr>
 213       </table>
 214 </div><div class="memdoc">
 215 <b>Value:</b><div class="fragment"><div class="line">({                                                                   \</div><div class="line">        acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0);                  \</div><div class="line">        acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0);                  \</div><div class="line">        acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0);                  \</div><div class="line">        acc.s1 = fma(src0.s2, weights_row0.s0, acc.s1);                  \</div><div class="line">        acc.s1 = fma(src0.s3, weights_row0.s1, acc.s1);                  \</div><div class="line">        acc.s1 = fma(src1.s0, weights_row0.s2, acc.s1);                  \</div><div class="line">    })</div></div><!-- fragment -->
 216 <p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00221">221</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
 217
 218 <p>Referenced by <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00360">depthwise_convolution_3x3_stridex2_stridey2_bifrost()</a>.</p>
 219
 220 </div>
 221 </div>
 222 <h2 class="groupheader">Function Documentation</h2>
 223 <a class="anchor" id="a2fbbf5b73260b66b310d4d857db5289e"></a>
 224 <div class="memitem">
 225 <div class="memproto">
 226       <table class="memname">
 227         <tr>
 228           <td class="memname">__kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost </td>
 229           <td>(</td>
 230           <td class="paramtype">__global uchar *&#160;</td>
 231           <td class="paramname"><em>src_ptr</em>, </td>
 232         </tr>
 233         <tr>
 234           <td class="paramkey"></td>
 235           <td></td>
 236           <td class="paramtype">uint&#160;</td>
 237           <td class="paramname"><em>src_stride_x</em>, </td>
 238         </tr>
 239         <tr>
 240           <td class="paramkey"></td>
 241           <td></td>
 242           <td class="paramtype">uint&#160;</td>
 243           <td class="paramname"><em>src_step_x</em>, </td>
 244         </tr>
 245         <tr>
 246           <td class="paramkey"></td>
 247           <td></td>
 248           <td class="paramtype">uint&#160;</td>
 249           <td class="paramname"><em>src_stride_y</em>, </td>
 250         </tr>
 251         <tr>
 252           <td class="paramkey"></td>
 253           <td></td>
 254           <td class="paramtype">uint&#160;</td>
 255           <td class="paramname"><em>src_step_y</em>, </td>
 256         </tr>
 257         <tr>
 258           <td class="paramkey"></td>
 259           <td></td>
 260           <td class="paramtype">uint&#160;</td>
 261           <td class="paramname"><em>src_stride_z</em>, </td>
 262         </tr>
 263         <tr>
 264           <td class="paramkey"></td>
 265           <td></td>
 266           <td class="paramtype">uint&#160;</td>
 267           <td class="paramname"><em>src_step_z</em>, </td>
 268         </tr>
 269         <tr>
 270           <td class="paramkey"></td>
 271           <td></td>
 272           <td class="paramtype">uint&#160;</td>
 273           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 274         </tr>
 275         <tr>
 276           <td class="paramkey"></td>
 277           <td></td>
 278           <td class="paramtype">__global uchar *&#160;</td>
 279           <td class="paramname"><em>dst_ptr</em>, </td>
 280         </tr>
 281         <tr>
 282           <td class="paramkey"></td>
 283           <td></td>
 284           <td class="paramtype">uint&#160;</td>
 285           <td class="paramname"><em>dst_stride_x</em>, </td>
 286         </tr>
 287         <tr>
 288           <td class="paramkey"></td>
 289           <td></td>
 290           <td class="paramtype">uint&#160;</td>
 291           <td class="paramname"><em>dst_step_x</em>, </td>
 292         </tr>
 293         <tr>
 294           <td class="paramkey"></td>
 295           <td></td>
 296           <td class="paramtype">uint&#160;</td>
 297           <td class="paramname"><em>dst_stride_y</em>, </td>
 298         </tr>
 299         <tr>
 300           <td class="paramkey"></td>
 301           <td></td>
 302           <td class="paramtype">uint&#160;</td>
 303           <td class="paramname"><em>dst_step_y</em>, </td>
 304         </tr>
 305         <tr>
 306           <td class="paramkey"></td>
 307           <td></td>
 308           <td class="paramtype">uint&#160;</td>
 309           <td class="paramname"><em>dst_stride_z</em>, </td>
 310         </tr>
 311         <tr>
 312           <td class="paramkey"></td>
 313           <td></td>
 314           <td class="paramtype">uint&#160;</td>
 315           <td class="paramname"><em>dst_step_z</em>, </td>
 316         </tr>
 317         <tr>
 318           <td class="paramkey"></td>
 319           <td></td>
 320           <td class="paramtype">uint&#160;</td>
 321           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
 322         </tr>
 323         <tr>
 324           <td class="paramkey"></td>
 325           <td></td>
 326           <td class="paramtype">__global uchar *&#160;</td>
 327           <td class="paramname"><em>weights_ptr</em>, </td>
 328         </tr>
 329         <tr>
 330           <td class="paramkey"></td>
 331           <td></td>
 332           <td class="paramtype">uint&#160;</td>
 333           <td class="paramname"><em>weights_stride_x</em>, </td>
 334         </tr>
 335         <tr>
 336           <td class="paramkey"></td>
 337           <td></td>
 338           <td class="paramtype">uint&#160;</td>
 339           <td class="paramname"><em>weights_step_x</em>, </td>
 340         </tr>
 341         <tr>
 342           <td class="paramkey"></td>
 343           <td></td>
 344           <td class="paramtype">uint&#160;</td>
 345           <td class="paramname"><em>weights_stride_y</em>, </td>
 346         </tr>
 347         <tr>
 348           <td class="paramkey"></td>
 349           <td></td>
 350           <td class="paramtype">uint&#160;</td>
 351           <td class="paramname"><em>weights_step_y</em>, </td>
 352         </tr>
 353         <tr>
 354           <td class="paramkey"></td>
 355           <td></td>
 356           <td class="paramtype">uint&#160;</td>
 357           <td class="paramname"><em>weights_stride_z</em>, </td>
 358         </tr>
 359         <tr>
 360           <td class="paramkey"></td>
 361           <td></td>
 362           <td class="paramtype">uint&#160;</td>
 363           <td class="paramname"><em>weights_step_z</em>, </td>
 364         </tr>
 365         <tr>
 366           <td class="paramkey"></td>
 367           <td></td>
 368           <td class="paramtype">uint&#160;</td>
 369           <td class="paramname"><em>weights_offset_first_element_in_bytes</em>, </td>
 370         </tr>
 371         <tr>
 372           <td class="paramkey"></td>
 373           <td></td>
 374           <td class="paramtype">__global uchar *&#160;</td>
 375           <td class="paramname"><em>biases_ptr</em>, </td>
 376         </tr>
 377         <tr>
 378           <td class="paramkey"></td>
 379           <td></td>
 380           <td class="paramtype">uint&#160;</td>
 381           <td class="paramname"><em>biases_stride_x</em>, </td>
 382         </tr>
 383         <tr>
 384           <td class="paramkey"></td>
 385           <td></td>
 386           <td class="paramtype">uint&#160;</td>
 387           <td class="paramname"><em>biases_step_x</em>, </td>
 388         </tr>
 389         <tr>
 390           <td class="paramkey"></td>
 391           <td></td>
 392           <td class="paramtype">uint&#160;</td>
 393           <td class="paramname"><em>biases_offset_first_element_in_bytes</em>&#160;</td>
 394         </tr>
 395         <tr>
 396           <td></td>
 397           <td>)</td>
 398           <td></td><td></td>
 399         </tr>
 400       </table>
 401 </div><div class="memdoc">
 402
 403 <p>This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 1. </p>
 404 <dl class="params"><dt>Parameters</dt><dd>
 405   <table class="params">
 406     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source image. Supported data types: F32 </td></tr>
 407     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
 408     <tr><td class="paramdir">[in]</td><td class="paramname">src_step_x</td><td>src_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 409     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
 410     <tr><td class="paramdir">[in]</td><td class="paramname">src_step_y</td><td>src_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 411     <tr><td class="paramdir">[in]</td><td class="paramname">src_offset_first_element_in_bytes</td><td>The offset of the first element in the source image </td></tr>
 412     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
 413     <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
 414     <tr><td class="paramdir">[in]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F32 </td></tr>
 415     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination tensor in X dimension (in bytes) </td></tr>
 416     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 417     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination tensor in Y dimension (in bytes) </td></tr>
 418     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 419     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
 420     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
 421     <tr><td class="paramdir">[in]</td><td class="paramname">dst_offset_first_element_in_bytes</td><td>The offset of the first element in the destination tensor </td></tr>
 422     <tr><td class="paramdir">[in]</td><td class="paramname">weights_ptr</td><td>Pointer to the weights tensor. Supported data types: F32 </td></tr>
 423     <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_x</td><td>Stride of the weights tensor in X dimension (in bytes) </td></tr>
 424     <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_x</td><td>weights_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 425     <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_y</td><td>Stride of the weights tensor in Y dimension (in bytes) </td></tr>
 426     <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_y</td><td>weights_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 427     <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_z</td><td>Stride of the weights tensor in Z dimension (in bytes) </td></tr>
 428     <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_z</td><td>weights_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
 429     <tr><td class="paramdir">[in]</td><td class="paramname">weights_offset_first_element_in_bytes</td><td>The offset of the first element in the biases vector </td></tr>
 430     <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>(Optional) Pointer to the biases vector. Supported data types: F32 </td></tr>
 431     <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>(Optional) Stride of the biases vector in X dimension (in bytes) </td></tr>
 432     <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>(Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 433     <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>(Optional) The offset of the first element in the biases vector </td></tr>
 434   </table>
 435   </dd>
 436 </dl>
 437
 438 <p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00263">263</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
 439
 440 <p>References <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00117">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00099">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>, <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00211">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00151">Tensor3D::ptr</a>, and <a class="el" href="helpers_8h_source.xhtml#l00290">vector_offset()</a>.</p>
 441 <div class="fragment"><div class="line"><a name="l00272"></a><span class="lineno">  272</span>&#160;{</div><div class="line"><a name="l00273"></a><span class="lineno">  273</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>    <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a>     = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00274"></a><span class="lineno">  274</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>    <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>     = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00275"></a><span class="lineno">  275</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> weights = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(weights);</div><div class="line"><a name="l00276"></a><span class="lineno">  276</span>&#160;</div><div class="line"><a name="l00277"></a><span class="lineno">  277</span>&#160;    float2 pixels0 = 0.0f;</div><div class="line"><a name="l00278"></a><span class="lineno">  278</span>&#160;    float2 pixels1 = 0.0f;</div><div class="line"><a name="l00279"></a><span class="lineno">  279</span>&#160;    float2 pixels2 = 0.0f;</div><div class="line"><a name="l00280"></a><span class="lineno">  280</span>&#160;    float2 pixels3 = 0.0f;</div><div class="line"><a name="l00281"></a><span class="lineno">  281</span>&#160;</div><div class="line"><a name="l00282"></a><span class="lineno">  282</span>&#160;    __global uchar *weights_addr = (__global uchar *)weights.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>;</div><div class="line"><a name="l00283"></a><span class="lineno">  283</span>&#160;    __global uchar *src_addr     = (__global uchar *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0);</div><div class="line"><a name="l00284"></a><span class="lineno">  284</span>&#160;</div><div class="line"><a name="l00285"></a><span class="lineno">  285</span>&#160;    <span class="comment">// Load the weights</span></div><div class="line"><a name="l00286"></a><span class="lineno">  286</span>&#160;    float3 weights_row0 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 0 * weights_stride_y));</div><div class="line"><a name="l00287"></a><span class="lineno">  287</span>&#160;    float3 weights_row1 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 1 * weights_stride_y));</div><div class="line"><a name="l00288"></a><span class="lineno">  288</span>&#160;    float3 weights_row2 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 2 * weights_stride_y));</div><div class="line"><a name="l00289"></a><span class="lineno">  289</span>&#160;</div><div class="line"><a name="l00290"></a><span class="lineno">  290</span>&#160;    <span class="comment">// Note: Since each work-item computes 4x2 elements, we need to load 4 rows from the input tensor</span></div><div class="line"><a name="l00291"></a><span class="lineno">  291</span>&#160;    float4 src00 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00292"></a><span class="lineno">  292</span>&#160;    float4 src10 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00293"></a><span class="lineno">  293</span>&#160;    float4 src20 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00294"></a><span class="lineno">  294</span>&#160;    float4 src30 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00295"></a><span class="lineno">  295</span>&#160;    float4 src40 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00296"></a><span class="lineno">  296</span>&#160;    float4 src50 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 5 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00297"></a><span class="lineno">  297</span>&#160;</div><div class="line"><a name="l00298"></a><span class="lineno">  298</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src00, weights_row0);</div><div class="line"><a name="l00299"></a><span class="lineno">  299</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src10, weights_row1);</div><div class="line"><a name="l00300"></a><span class="lineno">  300</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src20, weights_row2);</div><div class="line"><a name="l00301"></a><span class="lineno">  301</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src10, weights_row0);</div><div class="line"><a name="l00302"></a><span class="lineno">  302</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src20, weights_row1);</div><div class="line"><a name="l00303"></a><span class="lineno">  303</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src30, weights_row2);</div><div class="line"><a name="l00304"></a><span class="lineno">  304</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src20, weights_row0);</div><div class="line"><a name="l00305"></a><span class="lineno">  305</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src30, weights_row1);</div><div class="line"><a name="l00306"></a><span class="lineno">  306</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src40, weights_row2);</div><div class="line"><a name="l00307"></a><span class="lineno">  307</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src30, weights_row0);</div><div class="line"><a name="l00308"></a><span class="lineno">  308</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src40, weights_row1);</div><div class="line"><a name="l00309"></a><span class="lineno">  309</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src50, weights_row2);</div><div class="line"><a name="l00310"></a><span class="lineno">  310</span>&#160;</div><div class="line"><a name="l00311"></a><span class="lineno">  311</span>&#160;<span class="preprocessor">#ifdef HAS_BIAS</span></div><div class="line"><a name="l00312"></a><span class="lineno">  312</span>&#160;    <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>(biases);</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="keywordtype">float</span> bias = *((__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a>(&amp;biases, get_global_id(2))));</div><div class="line"><a name="l00315"></a><span class="lineno">  315</span>&#160;</div><div class="line"><a name="l00316"></a><span class="lineno">  316</span>&#160;    pixels0 += (float2)bias;</div><div class="line"><a name="l00317"></a><span class="lineno">  317</span>&#160;    pixels1 += (float2)bias;</div><div class="line"><a name="l00318"></a><span class="lineno">  318</span>&#160;    pixels2 += (float2)bias;</div><div class="line"><a name="l00319"></a><span class="lineno">  319</span>&#160;    pixels3 += (float2)bias;</div><div class="line"><a name="l00320"></a><span class="lineno">  320</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(HAS_BIAS) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00321"></a><span class="lineno">  321</span>&#160;</div><div class="line"><a name="l00322"></a><span class="lineno">  322</span>&#160;    vstore2(pixels0, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 0 * dst_stride_y));</div><div class="line"><a name="l00323"></a><span class="lineno">  323</span>&#160;    vstore2(pixels1, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 1 * dst_stride_y));</div><div class="line"><a name="l00324"></a><span class="lineno">  324</span>&#160;    vstore2(pixels2, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 2 * dst_stride_y));</div><div class="line"><a name="l00325"></a><span class="lineno">  325</span>&#160;    vstore2(pixels3, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 3 * dst_stride_y));</div><div class="line"><a name="l00326"></a><span class="lineno">  326</span>&#160;}</div><div class="ttc" id="struct_vector_xhtml"><div class="ttname"><a href="struct_vector.xhtml">Vector</a></div><div class="ttdoc">Structure to hold Vector information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00132">helpers.h:132</a></div></div>
 442 <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#l00117">helpers.h:117</a></div></div>
 443 <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#l00149">helpers.h:149</a></div></div>
 444 <div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
 445 <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_convolution_8cpp_source.xhtml#l00137">Convolution.cpp:137</a></div></div>
 446 <div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
 447 <div class="ttc" id="depthwise__convolution_8cl_xhtml_aa5ab4cf85fe56de3821727cfcef327d9"><div class="ttname"><a href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a></div><div class="ttdeci">#define CONVOLUTION1x3_BIFROST2X1_STRIDE1(acc, src0, weights_row0)</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution_8cl_source.xhtml#l00211">depthwise_convolution.cl:211</a></div></div>
 448 <div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::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#l00142">helpers.h:142</a></div></div>
 449 <div class="ttc" id="helpers_8h_xhtml_a64d779f80eeb923e0ab2313433f7b40b"><div class="ttname"><a href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
 450 <div class="ttc" id="helpers_8h_xhtml_a7e4940407322d6f0ccb8b6b86b856019"><div class="ttname"><a href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a></div><div class="ttdeci">__global const uchar * vector_offset(const Vector *vec, int x)</div><div class="ttdoc">Get the pointer position of a Vector. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00290">helpers.h:290</a></div></div>
 451 <div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
 452 <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#l00151">helpers.h:151</a></div></div>
 453 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
 454 </div><!-- fragment -->
 455 </div>
 456 </div>
 457 <a class="anchor" id="a328ab548c536089e8238fb33de617bec"></a>
 458 <div class="memitem">
 459 <div class="memproto">
 460       <table class="memname">
 461         <tr>
 462           <td class="memname">__kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost </td>
 463           <td>(</td>
 464           <td class="paramtype">__global uchar *&#160;</td>
 465           <td class="paramname"><em>src_ptr</em>, </td>
 466         </tr>
 467         <tr>
 468           <td class="paramkey"></td>
 469           <td></td>
 470           <td class="paramtype">uint&#160;</td>
 471           <td class="paramname"><em>src_stride_x</em>, </td>
 472         </tr>
 473         <tr>
 474           <td class="paramkey"></td>
 475           <td></td>
 476           <td class="paramtype">uint&#160;</td>
 477           <td class="paramname"><em>src_step_x</em>, </td>
 478         </tr>
 479         <tr>
 480           <td class="paramkey"></td>
 481           <td></td>
 482           <td class="paramtype">uint&#160;</td>
 483           <td class="paramname"><em>src_stride_y</em>, </td>
 484         </tr>
 485         <tr>
 486           <td class="paramkey"></td>
 487           <td></td>
 488           <td class="paramtype">uint&#160;</td>
 489           <td class="paramname"><em>src_step_y</em>, </td>
 490         </tr>
 491         <tr>
 492           <td class="paramkey"></td>
 493           <td></td>
 494           <td class="paramtype">uint&#160;</td>
 495           <td class="paramname"><em>src_stride_z</em>, </td>
 496         </tr>
 497         <tr>
 498           <td class="paramkey"></td>
 499           <td></td>
 500           <td class="paramtype">uint&#160;</td>
 501           <td class="paramname"><em>src_step_z</em>, </td>
 502         </tr>
 503         <tr>
 504           <td class="paramkey"></td>
 505           <td></td>
 506           <td class="paramtype">uint&#160;</td>
 507           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 508         </tr>
 509         <tr>
 510           <td class="paramkey"></td>
 511           <td></td>
 512           <td class="paramtype">__global uchar *&#160;</td>
 513           <td class="paramname"><em>dst_ptr</em>, </td>
 514         </tr>
 515         <tr>
 516           <td class="paramkey"></td>
 517           <td></td>
 518           <td class="paramtype">uint&#160;</td>
 519           <td class="paramname"><em>dst_stride_x</em>, </td>
 520         </tr>
 521         <tr>
 522           <td class="paramkey"></td>
 523           <td></td>
 524           <td class="paramtype">uint&#160;</td>
 525           <td class="paramname"><em>dst_step_x</em>, </td>
 526         </tr>
 527         <tr>
 528           <td class="paramkey"></td>
 529           <td></td>
 530           <td class="paramtype">uint&#160;</td>
 531           <td class="paramname"><em>dst_stride_y</em>, </td>
 532         </tr>
 533         <tr>
 534           <td class="paramkey"></td>
 535           <td></td>
 536           <td class="paramtype">uint&#160;</td>
 537           <td class="paramname"><em>dst_step_y</em>, </td>
 538         </tr>
 539         <tr>
 540           <td class="paramkey"></td>
 541           <td></td>
 542           <td class="paramtype">uint&#160;</td>
 543           <td class="paramname"><em>dst_stride_z</em>, </td>
 544         </tr>
 545         <tr>
 546           <td class="paramkey"></td>
 547           <td></td>
 548           <td class="paramtype">uint&#160;</td>
 549           <td class="paramname"><em>dst_step_z</em>, </td>
 550         </tr>
 551         <tr>
 552           <td class="paramkey"></td>
 553           <td></td>
 554           <td class="paramtype">uint&#160;</td>
 555           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
 556         </tr>
 557         <tr>
 558           <td class="paramkey"></td>
 559           <td></td>
 560           <td class="paramtype">__global uchar *&#160;</td>
 561           <td class="paramname"><em>weights_ptr</em>, </td>
 562         </tr>
 563         <tr>
 564           <td class="paramkey"></td>
 565           <td></td>
 566           <td class="paramtype">uint&#160;</td>
 567           <td class="paramname"><em>weights_stride_x</em>, </td>
 568         </tr>
 569         <tr>
 570           <td class="paramkey"></td>
 571           <td></td>
 572           <td class="paramtype">uint&#160;</td>
 573           <td class="paramname"><em>weights_step_x</em>, </td>
 574         </tr>
 575         <tr>
 576           <td class="paramkey"></td>
 577           <td></td>
 578           <td class="paramtype">uint&#160;</td>
 579           <td class="paramname"><em>weights_stride_y</em>, </td>
 580         </tr>
 581         <tr>
 582           <td class="paramkey"></td>
 583           <td></td>
 584           <td class="paramtype">uint&#160;</td>
 585           <td class="paramname"><em>weights_step_y</em>, </td>
 586         </tr>
 587         <tr>
 588           <td class="paramkey"></td>
 589           <td></td>
 590           <td class="paramtype">uint&#160;</td>
 591           <td class="paramname"><em>weights_stride_z</em>, </td>
 592         </tr>
 593         <tr>
 594           <td class="paramkey"></td>
 595           <td></td>
 596           <td class="paramtype">uint&#160;</td>
 597           <td class="paramname"><em>weights_step_z</em>, </td>
 598         </tr>
 599         <tr>
 600           <td class="paramkey"></td>
 601           <td></td>
 602           <td class="paramtype">uint&#160;</td>
 603           <td class="paramname"><em>weights_offset_first_element_in_bytes</em>, </td>
 604         </tr>
 605         <tr>
 606           <td class="paramkey"></td>
 607           <td></td>
 608           <td class="paramtype">__global uchar *&#160;</td>
 609           <td class="paramname"><em>biases_ptr</em>, </td>
 610         </tr>
 611         <tr>
 612           <td class="paramkey"></td>
 613           <td></td>
 614           <td class="paramtype">uint&#160;</td>
 615           <td class="paramname"><em>biases_stride_x</em>, </td>
 616         </tr>
 617         <tr>
 618           <td class="paramkey"></td>
 619           <td></td>
 620           <td class="paramtype">uint&#160;</td>
 621           <td class="paramname"><em>biases_step_x</em>, </td>
 622         </tr>
 623         <tr>
 624           <td class="paramkey"></td>
 625           <td></td>
 626           <td class="paramtype">uint&#160;</td>
 627           <td class="paramname"><em>biases_offset_first_element_in_bytes</em>&#160;</td>
 628         </tr>
 629         <tr>
 630           <td></td>
 631           <td>)</td>
 632           <td></td><td></td>
 633         </tr>
 634       </table>
 635 </div><div class="memdoc">
 636
 637 <p>This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 2. </p>
 638 <dl class="params"><dt>Parameters</dt><dd>
 639   <table class="params">
 640     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source image. Supported data types: F32 </td></tr>
 641     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
 642     <tr><td class="paramdir">[in]</td><td class="paramname">src_step_x</td><td>src_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 643     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
 644     <tr><td class="paramdir">[in]</td><td class="paramname">src_step_y</td><td>src_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 645     <tr><td class="paramdir">[in]</td><td class="paramname">src_offset_first_element_in_bytes</td><td>The offset of the first element in the source image </td></tr>
 646     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
 647     <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
 648     <tr><td class="paramdir">[in]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F32 </td></tr>
 649     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination tensor in X dimension (in bytes) </td></tr>
 650     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 651     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination tensor in Y dimension (in bytes) </td></tr>
 652     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 653     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
 654     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
 655     <tr><td class="paramdir">[in]</td><td class="paramname">dst_offset_first_element_in_bytes</td><td>The offset of the first element in the destination tensor </td></tr>
 656     <tr><td class="paramdir">[in]</td><td class="paramname">weights_ptr</td><td>Pointer to the weights tensor. Supported data types: F32 </td></tr>
 657     <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_x</td><td>Stride of the weights tensor in X dimension (in bytes) </td></tr>
 658     <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_x</td><td>weights_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 659     <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_y</td><td>Stride of the weights tensor in Y dimension (in bytes) </td></tr>
 660     <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_y</td><td>weights_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 661     <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_z</td><td>Stride of the weights tensor in Z dimension (in bytes) </td></tr>
 662     <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_z</td><td>weights_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
 663     <tr><td class="paramdir">[in]</td><td class="paramname">weights_offset_first_element_in_bytes</td><td>The offset of the first element in the biases vector </td></tr>
 664     <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>(Optional) Pointer to the biases vector. Supported data types: F32 </td></tr>
 665     <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>(Optional) Stride of the biases vector in X dimension (in bytes) </td></tr>
 666     <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>(Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 667     <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>(Optional) The offset of the first element in the biases vector </td></tr>
 668   </table>
 669   </dd>
 670 </dl>
 671
 672 <p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00360">360</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
 673
 674 <p>References <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00117">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00096">CONVERT_TO_VECTOR_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00099">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>, <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00221">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00066">IMAGE_DECLARATION</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00134">Vector::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00151">Tensor3D::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00074">TENSOR3D_DECLARATION</a>, <a class="el" href="helpers_8h_source.xhtml#l00060">VECTOR_DECLARATION</a>, and <a class="el" href="helpers_8h_source.xhtml#l00290">vector_offset()</a>.</p>
 675 <div class="fragment"><div class="line"><a name="l00369"></a><span class="lineno">  369</span>&#160;{</div><div class="line"><a name="l00370"></a><span class="lineno">  370</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>    <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a>     = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00371"></a><span class="lineno">  371</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>    <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>     = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00372"></a><span class="lineno">  372</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> weights = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(weights);</div><div class="line"><a name="l00373"></a><span class="lineno">  373</span>&#160;</div><div class="line"><a name="l00374"></a><span class="lineno">  374</span>&#160;    float2 pixels0 = 0.0f;</div><div class="line"><a name="l00375"></a><span class="lineno">  375</span>&#160;    float2 pixels1 = 0.0f;</div><div class="line"><a name="l00376"></a><span class="lineno">  376</span>&#160;</div><div class="line"><a name="l00377"></a><span class="lineno">  377</span>&#160;    __global uchar *weights_addr = (__global uchar *)weights.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>;</div><div class="line"><a name="l00378"></a><span class="lineno">  378</span>&#160;    __global uchar *src_addr     = (__global uchar *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0);</div><div class="line"><a name="l00379"></a><span class="lineno">  379</span>&#160;</div><div class="line"><a name="l00380"></a><span class="lineno">  380</span>&#160;    <span class="comment">// Load the weights</span></div><div class="line"><a name="l00381"></a><span class="lineno">  381</span>&#160;    float3 weights_row0 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 0 * weights_stride_y));</div><div class="line"><a name="l00382"></a><span class="lineno">  382</span>&#160;    float3 weights_row1 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 1 * weights_stride_y));</div><div class="line"><a name="l00383"></a><span class="lineno">  383</span>&#160;    float3 weights_row2 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 2 * weights_stride_y));</div><div class="line"><a name="l00384"></a><span class="lineno">  384</span>&#160;</div><div class="line"><a name="l00385"></a><span class="lineno">  385</span>&#160;    <span class="comment">// Note: Since each work-item computes 4x2 elements, we need to load 5 rows from the input tensor</span></div><div class="line"><a name="l00386"></a><span class="lineno">  386</span>&#160;    float4 src00 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00387"></a><span class="lineno">  387</span>&#160;    float2 src01 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00388"></a><span class="lineno">  388</span>&#160;    float4 src10 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00389"></a><span class="lineno">  389</span>&#160;    float2 src11 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00390"></a><span class="lineno">  390</span>&#160;    float4 src20 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00391"></a><span class="lineno">  391</span>&#160;    float2 src21 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00392"></a><span class="lineno">  392</span>&#160;    float4 src30 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00393"></a><span class="lineno">  393</span>&#160;    float2 src31 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00394"></a><span class="lineno">  394</span>&#160;    float4 src40 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row4</span></div><div class="line"><a name="l00395"></a><span class="lineno">  395</span>&#160;    float2 src41 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row4</span></div><div class="line"><a name="l00396"></a><span class="lineno">  396</span>&#160;</div><div class="line"><a name="l00397"></a><span class="lineno">  397</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src00, src01, weights_row0);</div><div class="line"><a name="l00398"></a><span class="lineno">  398</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src10, src11, weights_row1);</div><div class="line"><a name="l00399"></a><span class="lineno">  399</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src20, src21, weights_row2);</div><div class="line"><a name="l00400"></a><span class="lineno">  400</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src20, src21, weights_row0);</div><div class="line"><a name="l00401"></a><span class="lineno">  401</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src30, src31, weights_row1);</div><div class="line"><a name="l00402"></a><span class="lineno">  402</span>&#160;    <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src40, src41, weights_row2);</div><div class="line"><a name="l00403"></a><span class="lineno">  403</span>&#160;</div><div class="line"><a name="l00404"></a><span class="lineno">  404</span>&#160;<span class="preprocessor">#ifdef HAS_BIAS</span></div><div class="line"><a name="l00405"></a><span class="lineno">  405</span>&#160;    <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>(biases);</div><div class="line"><a name="l00406"></a><span class="lineno">  406</span>&#160;</div><div class="line"><a name="l00407"></a><span class="lineno">  407</span>&#160;    <span class="keywordtype">float</span> bias = *((__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a>(&amp;biases, get_global_id(2))));</div><div class="line"><a name="l00408"></a><span class="lineno">  408</span>&#160;</div><div class="line"><a name="l00409"></a><span class="lineno">  409</span>&#160;    pixels0 += (float2)bias;</div><div class="line"><a name="l00410"></a><span class="lineno">  410</span>&#160;    pixels1 += (float2)bias;</div><div class="line"><a name="l00411"></a><span class="lineno">  411</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(HAS_BIAS) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00412"></a><span class="lineno">  412</span>&#160;</div><div class="line"><a name="l00413"></a><span class="lineno">  413</span>&#160;    vstore2(pixels0, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 0 * dst_stride_y));</div><div class="line"><a name="l00414"></a><span class="lineno">  414</span>&#160;    vstore2(pixels1, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 1 * dst_stride_y));</div><div class="line"><a name="l00415"></a><span class="lineno">  415</span>&#160;}</div><div class="ttc" id="struct_vector_xhtml"><div class="ttname"><a href="struct_vector.xhtml">Vector</a></div><div class="ttdoc">Structure to hold Vector information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00132">helpers.h:132</a></div></div>
 676 <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#l00117">helpers.h:117</a></div></div>
 677 <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#l00149">helpers.h:149</a></div></div>
 678 <div class="ttc" id="depthwise__convolution_8cl_xhtml_a31cbdd6b28ee6331839432f69d87fdcb"><div class="ttname"><a href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a></div><div class="ttdeci">#define CONVOLUTION1x3_BIFROST2X1_STRIDE2(acc, src0, src1, weights_row0)</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution_8cl_source.xhtml#l00221">depthwise_convolution.cl:221</a></div></div>
 679 <div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
 680 <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_convolution_8cpp_source.xhtml#l00137">Convolution.cpp:137</a></div></div>
 681 <div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
 682 <div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::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#l00142">helpers.h:142</a></div></div>
 683 <div class="ttc" id="helpers_8h_xhtml_a64d779f80eeb923e0ab2313433f7b40b"><div class="ttname"><a href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
 684 <div class="ttc" id="helpers_8h_xhtml_a7e4940407322d6f0ccb8b6b86b856019"><div class="ttname"><a href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a></div><div class="ttdeci">__global const uchar * vector_offset(const Vector *vec, int x)</div><div class="ttdoc">Get the pointer position of a Vector. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00290">helpers.h:290</a></div></div>
 685 <div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
 686 <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#l00151">helpers.h:151</a></div></div>
 687 <div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
 688 </div><!-- fragment -->
 689 </div>
 690 </div>
 691 </div><!-- contents -->
 692 </div><!-- doc-content -->
 693 <!-- start footer part -->
 694 <div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
 695   <ul>
 696     <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="depthwise__convolution_8cl.xhtml">depthwise_convolution.cl</a></li>
 697     <li class="footer">Generated on Thu Feb 22 2018 15:45:22 for Compute Library by
 698     <a href="http://www.doxygen.org/index.html">
 699     <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
 700   </ul>
 701 </div>
 702 </body>
 703 </html>