documentation/convolution__layer_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>ARM Compute Library: src/core/CL/cl_kernels/convolution_layer.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">ARM Compute Library
  43    &#160;<span id="projectnumber">17.04</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="namespaces.xhtml"><span>Namespaces</span></a></li>
  59       <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
  60       <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
  61       <li>
  62         <div id="MSearchBox" class="MSearchBoxInactive">
  63         <span class="left">
  64           <img id="MSearchSelect" src="search/mag_sel.png"
  65                onmouseover="return searchBox.OnSearchSelectShow()"
  66                onmouseout="return searchBox.OnSearchSelectHide()"
  67                alt=""/>
  68           <input type="text" id="MSearchField" value="Search" accesskey="S"
  69                onfocus="searchBox.OnSearchFieldFocus(true)"
  70                onblur="searchBox.OnSearchFieldFocus(false)"
  71                onkeyup="searchBox.OnSearchFieldChange(event)"/>
  72           </span><span class="right">
  73             <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
  74           </span>
  75         </div>
  76       </li>
  77     </ul>
  78   </div>
  79   <div id="navrow2" class="tabs2">
  80     <ul class="tablist">
  81       <li><a href="files.xhtml"><span>File&#160;List</span></a></li>
  82       <li><a href="globals.xhtml"><span>Globals</span></a></li>
  83     </ul>
  84   </div>
  85 </div><!-- top -->
  86 <div id="side-nav" class="ui-resizable side-nav-resizable">
  87   <div id="nav-tree">
  88     <div id="nav-tree-contents">
  89       <div id="nav-sync" class="sync"></div>
  90     </div>
  91   </div>
  92   <div id="splitbar" style="-moz-user-select:none;"
  93        class="ui-resizable-handle">
  94   </div>
  95 </div>
  96 <script type="text/javascript">
  97 $(document).ready(function(){initNavTree('convolution__layer_8cl.xhtml','');});
  98 </script>
  99 <div id="doc-content">
 100 <!-- window showing the filter options -->
 101 <div id="MSearchSelectWindow"
 102      onmouseover="return searchBox.OnSearchSelectShow()"
 103      onmouseout="return searchBox.OnSearchSelectHide()"
 104      onkeydown="return searchBox.OnSearchSelectKey(event)">
 105 </div>
 106
 107 <!-- iframe showing the search results (closed by default) -->
 108 <div id="MSearchResultsWindow">
 109 <iframe src="javascript:void(0)" frameborder="0"
 110         name="MSearchResults" id="MSearchResults">
 111 </iframe>
 112 </div>
 113
 114 <div class="header">
 115   <div class="summary">
 116 <a href="#func-members">Functions</a>  </div>
 117   <div class="headertitle">
 118 <div class="title">convolution_layer.cl File Reference</div>  </div>
 119 </div><!--header-->
 120 <div class="contents">
 121 <div class="textblock"><code>#include &quot;<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
 122 </div>
 123 <p><a href="convolution__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
 124 <table class="memberdecls">
 125 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
 126 Functions</h2></td></tr>
 127 <tr class="memitem:a229c10cbfba415bce4cc33d4530bfbd0"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="convolution__layer_8cl.xhtml#a229c10cbfba415bce4cc33d4530bfbd0">reshape_to_columns</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_offset_first_element_in_bytes, __global uchar *bias_ptr, uint bias_stride_x, uint bias_step_x, uint bias_stride_y, uint bias_step_y, uint bias_offset_first_element_in_bytes, uint width, uint height)</td></tr>
 128 <tr class="memdesc:a229c10cbfba415bce4cc33d4530bfbd0"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel reshapes the tensor's low three dimensions to single column.  <a href="#a229c10cbfba415bce4cc33d4530bfbd0">More...</a><br /></td></tr>
 129 <tr class="separator:a229c10cbfba415bce4cc33d4530bfbd0"><td class="memSeparator" colspan="2">&#160;</td></tr>
 130 <tr class="memitem:afaddfccc926506181141b2dc86cb73f1"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="convolution__layer_8cl.xhtml#afaddfccc926506181141b2dc86cb73f1">im2col_generic</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_offset_first_element_in_bytes, int kernel_size, int kernel_depth, int width, int2 input_dims, int2 strides, int2 paddings)</td></tr>
 131 <tr class="memdesc:afaddfccc926506181141b2dc86cb73f1"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel performs a reshaping of the input tensor to a tensor used to perform convolution using GEMM.  <a href="#afaddfccc926506181141b2dc86cb73f1">More...</a><br /></td></tr>
 132 <tr class="separator:afaddfccc926506181141b2dc86cb73f1"><td class="memSeparator" colspan="2">&#160;</td></tr>
 133 <tr class="memitem:af11bad06b2f7c26cc6ee2f93add5dcbf"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="convolution__layer_8cl.xhtml#af11bad06b2f7c26cc6ee2f93add5dcbf">col2im</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, 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, uint width)</td></tr>
 134 <tr class="memdesc:af11bad06b2f7c26cc6ee2f93add5dcbf"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel performs a reshaping of the output of the convolution layer.  <a href="#af11bad06b2f7c26cc6ee2f93add5dcbf">More...</a><br /></td></tr>
 135 <tr class="separator:af11bad06b2f7c26cc6ee2f93add5dcbf"><td class="memSeparator" colspan="2">&#160;</td></tr>
 136 <tr class="memitem:a6f31f2c8d606e1b57d9bdfacd416024c"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="convolution__layer_8cl.xhtml#a6f31f2c8d606e1b57d9bdfacd416024c">im2col_reduced</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_offset_first_element_in_bytes, uint width, uint height)</td></tr>
 137 <tr class="memdesc:a6f31f2c8d606e1b57d9bdfacd416024c"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel reshapes the tensor's low three dimensions to single row for GEMM operation.  <a href="#a6f31f2c8d606e1b57d9bdfacd416024c">More...</a><br /></td></tr>
 138 <tr class="separator:a6f31f2c8d606e1b57d9bdfacd416024c"><td class="memSeparator" colspan="2">&#160;</td></tr>
 139 </table>
 140 <h2 class="groupheader">Function Documentation</h2>
 141 <a class="anchor" id="af11bad06b2f7c26cc6ee2f93add5dcbf"></a>
 142 <div class="memitem">
 143 <div class="memproto">
 144       <table class="memname">
 145         <tr>
 146           <td class="memname">__kernel void col2im </td>
 147           <td>(</td>
 148           <td class="paramtype">__global uchar *&#160;</td>
 149           <td class="paramname"><em>src_ptr</em>, </td>
 150         </tr>
 151         <tr>
 152           <td class="paramkey"></td>
 153           <td></td>
 154           <td class="paramtype">uint&#160;</td>
 155           <td class="paramname"><em>src_stride_x</em>, </td>
 156         </tr>
 157         <tr>
 158           <td class="paramkey"></td>
 159           <td></td>
 160           <td class="paramtype">uint&#160;</td>
 161           <td class="paramname"><em>src_step_x</em>, </td>
 162         </tr>
 163         <tr>
 164           <td class="paramkey"></td>
 165           <td></td>
 166           <td class="paramtype">uint&#160;</td>
 167           <td class="paramname"><em>src_stride_y</em>, </td>
 168         </tr>
 169         <tr>
 170           <td class="paramkey"></td>
 171           <td></td>
 172           <td class="paramtype">uint&#160;</td>
 173           <td class="paramname"><em>src_step_y</em>, </td>
 174         </tr>
 175         <tr>
 176           <td class="paramkey"></td>
 177           <td></td>
 178           <td class="paramtype">uint&#160;</td>
 179           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 180         </tr>
 181         <tr>
 182           <td class="paramkey"></td>
 183           <td></td>
 184           <td class="paramtype">__global uchar *&#160;</td>
 185           <td class="paramname"><em>dst_ptr</em>, </td>
 186         </tr>
 187         <tr>
 188           <td class="paramkey"></td>
 189           <td></td>
 190           <td class="paramtype">uint&#160;</td>
 191           <td class="paramname"><em>dst_stride_x</em>, </td>
 192         </tr>
 193         <tr>
 194           <td class="paramkey"></td>
 195           <td></td>
 196           <td class="paramtype">uint&#160;</td>
 197           <td class="paramname"><em>dst_step_x</em>, </td>
 198         </tr>
 199         <tr>
 200           <td class="paramkey"></td>
 201           <td></td>
 202           <td class="paramtype">uint&#160;</td>
 203           <td class="paramname"><em>dst_stride_y</em>, </td>
 204         </tr>
 205         <tr>
 206           <td class="paramkey"></td>
 207           <td></td>
 208           <td class="paramtype">uint&#160;</td>
 209           <td class="paramname"><em>dst_step_y</em>, </td>
 210         </tr>
 211         <tr>
 212           <td class="paramkey"></td>
 213           <td></td>
 214           <td class="paramtype">uint&#160;</td>
 215           <td class="paramname"><em>dst_stride_z</em>, </td>
 216         </tr>
 217         <tr>
 218           <td class="paramkey"></td>
 219           <td></td>
 220           <td class="paramtype">uint&#160;</td>
 221           <td class="paramname"><em>dst_step_z</em>, </td>
 222         </tr>
 223         <tr>
 224           <td class="paramkey"></td>
 225           <td></td>
 226           <td class="paramtype">uint&#160;</td>
 227           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
 228         </tr>
 229         <tr>
 230           <td class="paramkey"></td>
 231           <td></td>
 232           <td class="paramtype">uint&#160;</td>
 233           <td class="paramname"><em>width</em>&#160;</td>
 234         </tr>
 235         <tr>
 236           <td></td>
 237           <td>)</td>
 238           <td></td><td></td>
 239         </tr>
 240       </table>
 241 </div><div class="memdoc">
 242
 243 <p>This kernel performs a reshaping of the output of the convolution layer. </p>
 244 <dl class="section note"><dt>Note</dt><dd>The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float</dd></dl>
 245 <dl class="params"><dt>Parameters</dt><dd>
 246   <table class="params">
 247     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: F16, F32 </td></tr>
 248     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source tensor in X dimension (in bytes) </td></tr>
 249     <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>
 250     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source tensor in Y dimension (in bytes) </td></tr>
 251     <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>
 252     <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 tensor </td></tr>
 253     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F16, F32 </td></tr>
 254     <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>
 255     <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>
 256     <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>
 257     <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>
 258     <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>
 259     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
 260     <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>
 261     <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>The output tensor width </td></tr>
 262   </table>
 263   </dd>
 264 </dl>
 265
 266 <p>Definition at line <a class="el" href="convolution__layer_8cl_source.xhtml#l00175">175</a> of file <a class="el" href="convolution__layer_8cl_source.xhtml">convolution_layer.cl</a>.</p>
 267
 268 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00083">CONVERT_TO_TENSOR3D_STRUCT_NO_STEP</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00106">Tensor3D::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00108">Tensor3D::stride_x</a>, <a class="el" href="helpers_8h_source.xhtml#l00109">Tensor3D::stride_y</a>, and <a class="el" href="helpers_8h_source.xhtml#l00110">Tensor3D::stride_z</a>.</p>
 269 <div class="fragment"><div class="line"><a name="l00179"></a><span class="lineno">  179</span>&#160;{</div><div class="line"><a name="l00180"></a><span class="lineno">  180</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>    src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00181"></a><span class="lineno">  181</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> dst = <a class="code" href="helpers_8h.xhtml#a79e8e562daa6599317d2d1cd86ef1bf2">CONVERT_TO_TENSOR3D_STRUCT_NO_STEP</a>(dst);</div><div class="line"><a name="l00182"></a><span class="lineno">  182</span>&#160;</div><div class="line"><a name="l00183"></a><span class="lineno">  183</span>&#160;    <span class="keywordtype">int</span>      idx                         = get_global_id(0) * dst.<a class="code" href="struct_tensor3_d.xhtml#ad5ff7a2b2bd0eec50fe09c254b127d1c">stride_z</a> + (get_global_id(1) / width) * dst.<a class="code" href="struct_tensor3_d.xhtml#a4f0b90c9ecd6e57ceb3f37332fefe8f1">stride_y</a> + (get_global_id(1) % width) * dst.<a class="code" href="struct_tensor3_d.xhtml#ae01febbfd0689ef709f3ff6fdd2abc7e">stride_x</a>;</div><div class="line"><a name="l00184"></a><span class="lineno">  184</span>&#160;    __global uchar *tmp_out_ptr          = dst.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + idx;</div><div class="line"><a name="l00185"></a><span class="lineno">  185</span>&#160;    *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)(src.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>));</div><div class="line"><a name="l00186"></a><span class="lineno">  186</span>&#160;}</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>
 270 <div class="ttc" id="struct_tensor3_d_xhtml_ad5ff7a2b2bd0eec50fe09c254b127d1c"><div class="ttname"><a href="struct_tensor3_d.xhtml#ad5ff7a2b2bd0eec50fe09c254b127d1c">Tensor3D::stride_z</a></div><div class="ttdeci">int stride_z</div><div class="ttdoc">Stride of the image in Z dimension (in bytes) </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00110">helpers.h:110</a></div></div>
 271 <div class="ttc" id="struct_tensor3_d_xhtml_ae01febbfd0689ef709f3ff6fdd2abc7e"><div class="ttname"><a href="struct_tensor3_d.xhtml#ae01febbfd0689ef709f3ff6fdd2abc7e">Tensor3D::stride_x</a></div><div class="ttdeci">int stride_x</div><div class="ttdoc">Stride of the image in X dimension (in bytes) </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00108">helpers.h:108</a></div></div>
 272 <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#l00104">helpers.h:104</a></div></div>
 273 <div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00073">helpers.h:73</a></div></div>
 274 <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#l00095">helpers.h:95</a></div></div>
 275 <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#l00097">helpers.h:97</a></div></div>
 276 <div class="ttc" id="helpers_8h_xhtml_a79e8e562daa6599317d2d1cd86ef1bf2"><div class="ttname"><a href="helpers_8h.xhtml#a79e8e562daa6599317d2d1cd86ef1bf2">CONVERT_TO_TENSOR3D_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00083">helpers.h:83</a></div></div>
 277 <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#l00106">helpers.h:106</a></div></div>
 278 <div class="ttc" id="struct_tensor3_d_xhtml_a4f0b90c9ecd6e57ceb3f37332fefe8f1"><div class="ttname"><a href="struct_tensor3_d.xhtml#a4f0b90c9ecd6e57ceb3f37332fefe8f1">Tensor3D::stride_y</a></div><div class="ttdeci">int stride_y</div><div class="ttdoc">Stride of the image in Y dimension (in bytes) </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00109">helpers.h:109</a></div></div>
 279 </div><!-- fragment -->
 280 </div>
 281 </div>
 282 <a class="anchor" id="afaddfccc926506181141b2dc86cb73f1"></a>
 283 <div class="memitem">
 284 <div class="memproto">
 285       <table class="memname">
 286         <tr>
 287           <td class="memname">__kernel void im2col_generic </td>
 288           <td>(</td>
 289           <td class="paramtype">__global uchar *&#160;</td>
 290           <td class="paramname"><em>src_ptr</em>, </td>
 291         </tr>
 292         <tr>
 293           <td class="paramkey"></td>
 294           <td></td>
 295           <td class="paramtype">uint&#160;</td>
 296           <td class="paramname"><em>src_stride_x</em>, </td>
 297         </tr>
 298         <tr>
 299           <td class="paramkey"></td>
 300           <td></td>
 301           <td class="paramtype">uint&#160;</td>
 302           <td class="paramname"><em>src_step_x</em>, </td>
 303         </tr>
 304         <tr>
 305           <td class="paramkey"></td>
 306           <td></td>
 307           <td class="paramtype">uint&#160;</td>
 308           <td class="paramname"><em>src_stride_y</em>, </td>
 309         </tr>
 310         <tr>
 311           <td class="paramkey"></td>
 312           <td></td>
 313           <td class="paramtype">uint&#160;</td>
 314           <td class="paramname"><em>src_step_y</em>, </td>
 315         </tr>
 316         <tr>
 317           <td class="paramkey"></td>
 318           <td></td>
 319           <td class="paramtype">uint&#160;</td>
 320           <td class="paramname"><em>src_stride_z</em>, </td>
 321         </tr>
 322         <tr>
 323           <td class="paramkey"></td>
 324           <td></td>
 325           <td class="paramtype">uint&#160;</td>
 326           <td class="paramname"><em>src_step_z</em>, </td>
 327         </tr>
 328         <tr>
 329           <td class="paramkey"></td>
 330           <td></td>
 331           <td class="paramtype">uint&#160;</td>
 332           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 333         </tr>
 334         <tr>
 335           <td class="paramkey"></td>
 336           <td></td>
 337           <td class="paramtype">__global uchar *&#160;</td>
 338           <td class="paramname"><em>dst_ptr</em>, </td>
 339         </tr>
 340         <tr>
 341           <td class="paramkey"></td>
 342           <td></td>
 343           <td class="paramtype">uint&#160;</td>
 344           <td class="paramname"><em>dst_stride_x</em>, </td>
 345         </tr>
 346         <tr>
 347           <td class="paramkey"></td>
 348           <td></td>
 349           <td class="paramtype">uint&#160;</td>
 350           <td class="paramname"><em>dst_step_x</em>, </td>
 351         </tr>
 352         <tr>
 353           <td class="paramkey"></td>
 354           <td></td>
 355           <td class="paramtype">uint&#160;</td>
 356           <td class="paramname"><em>dst_stride_y</em>, </td>
 357         </tr>
 358         <tr>
 359           <td class="paramkey"></td>
 360           <td></td>
 361           <td class="paramtype">uint&#160;</td>
 362           <td class="paramname"><em>dst_step_y</em>, </td>
 363         </tr>
 364         <tr>
 365           <td class="paramkey"></td>
 366           <td></td>
 367           <td class="paramtype">uint&#160;</td>
 368           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
 369         </tr>
 370         <tr>
 371           <td class="paramkey"></td>
 372           <td></td>
 373           <td class="paramtype">int&#160;</td>
 374           <td class="paramname"><em>kernel_size</em>, </td>
 375         </tr>
 376         <tr>
 377           <td class="paramkey"></td>
 378           <td></td>
 379           <td class="paramtype">int&#160;</td>
 380           <td class="paramname"><em>kernel_depth</em>, </td>
 381         </tr>
 382         <tr>
 383           <td class="paramkey"></td>
 384           <td></td>
 385           <td class="paramtype">int&#160;</td>
 386           <td class="paramname"><em>width</em>, </td>
 387         </tr>
 388         <tr>
 389           <td class="paramkey"></td>
 390           <td></td>
 391           <td class="paramtype">int2&#160;</td>
 392           <td class="paramname"><em>input_dims</em>, </td>
 393         </tr>
 394         <tr>
 395           <td class="paramkey"></td>
 396           <td></td>
 397           <td class="paramtype">int2&#160;</td>
 398           <td class="paramname"><em>strides</em>, </td>
 399         </tr>
 400         <tr>
 401           <td class="paramkey"></td>
 402           <td></td>
 403           <td class="paramtype">int2&#160;</td>
 404           <td class="paramname"><em>paddings</em>&#160;</td>
 405         </tr>
 406         <tr>
 407           <td></td>
 408           <td>)</td>
 409           <td></td><td></td>
 410         </tr>
 411       </table>
 412 </div><div class="memdoc">
 413
 414 <p>This kernel performs a reshaping of the input tensor to a tensor used to perform convolution using GEMM. </p>
 415 <dl class="section note"><dt>Note</dt><dd>The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float </dd>
 416 <dd>
 417 In case biases will be added to the convolution -DHAS_BIAS has to be passed to append the final matrix with 1 in each row.</dd></dl>
 418 <dl class="params"><dt>Parameters</dt><dd>
 419   <table class="params">
 420     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: F16, F32 </td></tr>
 421     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source tensor in X dimension (in bytes) </td></tr>
 422     <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>
 423     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source tensor in Y dimension (in bytes) </td></tr>
 424     <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>
 425     <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>
 426     <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
 427     <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 tensor </td></tr>
 428     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F16, F32 </td></tr>
 429     <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>
 430     <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>
 431     <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>
 432     <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>
 433     <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>
 434     <tr><td class="paramdir">[in]</td><td class="paramname">kernel_size</td><td>The convolution kernel size </td></tr>
 435     <tr><td class="paramdir">[in]</td><td class="paramname">kernel_depth</td><td>The kernel depth </td></tr>
 436     <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>The output tensor width </td></tr>
 437     <tr><td class="paramdir">[in]</td><td class="paramname">input_dims</td><td>The input tensor dimensions </td></tr>
 438     <tr><td class="paramdir">[in]</td><td class="paramname">strides</td><td>The strides of the im2col operation </td></tr>
 439     <tr><td class="paramdir">[in]</td><td class="paramname">paddings</td><td>The input tensor paddings </td></tr>
 440   </table>
 441   </dd>
 442 </dl>
 443
 444 <p>Definition at line <a class="el" href="convolution__layer_8cl_source.xhtml#l00105">105</a> of file <a class="el" href="convolution__layer_8cl_source.xhtml">convolution_layer.cl</a>.</p>
 445
 446 <p>References <a class="el" href="helpers_8h_source.xhtml#l00076">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>, <a class="el" href="helpers_8h_source.xhtml#l00079">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00099">Image::stride_x</a>, <a class="el" href="helpers_8h_source.xhtml#l00100">Image::stride_y</a>, and <a class="el" href="helpers_8h_source.xhtml#l00213">tensor3D_offset()</a>.</p>
 447 <div class="fragment"><div class="line"><a name="l00114"></a><span class="lineno">  114</span>&#160;{</div><div class="line"><a name="l00115"></a><span class="lineno">  115</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> src = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(src);</div><div class="line"><a name="l00116"></a><span class="lineno">  116</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>    dst = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(dst);</div><div class="line"><a name="l00117"></a><span class="lineno">  117</span>&#160;</div><div class="line"><a name="l00118"></a><span class="lineno">  118</span>&#160;    <span class="comment">// Kernel radius</span></div><div class="line"><a name="l00119"></a><span class="lineno">  119</span>&#160;    <span class="keywordtype">int</span> KH = kernel_size &gt;&gt; 1;</div><div class="line"><a name="l00120"></a><span class="lineno">  120</span>&#160;</div><div class="line"><a name="l00121"></a><span class="lineno">  121</span>&#160;    <span class="comment">// Determine output index</span></div><div class="line"><a name="l00122"></a><span class="lineno">  122</span>&#160;    uint     idx                = (get_global_id(1) * width + get_global_id(0)) * dst.<a class="code" href="struct_image.xhtml#a4f0b90c9ecd6e57ceb3f37332fefe8f1">stride_y</a>;</div><div class="line"><a name="l00123"></a><span class="lineno">  123</span>&#160;    __global uchar *tmp_out_ptr = dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + idx;</div><div class="line"><a name="l00124"></a><span class="lineno">  124</span>&#160;</div><div class="line"><a name="l00125"></a><span class="lineno">  125</span>&#160;    <span class="comment">// Determine current input index</span></div><div class="line"><a name="l00126"></a><span class="lineno">  126</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> wi = get_global_id(0) * strides.x - paddings.x;</div><div class="line"><a name="l00127"></a><span class="lineno">  127</span>&#160;    <span class="keyword">const</span> <span class="keywordtype">int</span> hi = get_global_id(1) * strides.y - paddings.y;</div><div class="line"><a name="l00128"></a><span class="lineno">  128</span>&#160;</div><div class="line"><a name="l00129"></a><span class="lineno">  129</span>&#160;    <span class="comment">// Linearize convolution elements</span></div><div class="line"><a name="l00130"></a><span class="lineno">  130</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">int</span> d = 0; d &lt; kernel_depth; ++d)</div><div class="line"><a name="l00131"></a><span class="lineno">  131</span>&#160;    {</div><div class="line"><a name="l00132"></a><span class="lineno">  132</span>&#160;        <span class="keywordtype">int</span> cur_h = hi;</div><div class="line"><a name="l00133"></a><span class="lineno">  133</span>&#160;        <span class="keywordflow">for</span>(<span class="keywordtype">int</span> h = -KH; h &lt;= KH; ++h, ++cur_h)</div><div class="line"><a name="l00134"></a><span class="lineno">  134</span>&#160;        {</div><div class="line"><a name="l00135"></a><span class="lineno">  135</span>&#160;            <span class="keywordtype">int</span> cur_w = wi;</div><div class="line"><a name="l00136"></a><span class="lineno">  136</span>&#160;            <span class="keywordflow">for</span>(<span class="keywordtype">int</span> w = -KH; w &lt;= KH; ++w, ++cur_w)</div><div class="line"><a name="l00137"></a><span class="lineno">  137</span>&#160;            {</div><div class="line"><a name="l00138"></a><span class="lineno">  138</span>&#160;                <span class="keywordflow">if</span>(((cur_w &lt; 0) || (cur_w &gt;= input_dims.x)) || ((cur_h &lt; 0) || (cur_h &gt;= input_dims.y)))</div><div class="line"><a name="l00139"></a><span class="lineno">  139</span>&#160;                {</div><div class="line"><a name="l00140"></a><span class="lineno">  140</span>&#160;                    *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = 0;</div><div class="line"><a name="l00141"></a><span class="lineno">  141</span>&#160;                }</div><div class="line"><a name="l00142"></a><span class="lineno">  142</span>&#160;                <span class="keywordflow">else</span></div><div class="line"><a name="l00143"></a><span class="lineno">  143</span>&#160;                {</div><div class="line"><a name="l00144"></a><span class="lineno">  144</span>&#160;                    *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)(<a class="code" href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;src, w, h, d)));</div><div class="line"><a name="l00145"></a><span class="lineno">  145</span>&#160;                }</div><div class="line"><a name="l00146"></a><span class="lineno">  146</span>&#160;                tmp_out_ptr += dst.<a class="code" href="struct_image.xhtml#ae01febbfd0689ef709f3ff6fdd2abc7e">stride_x</a>;</div><div class="line"><a name="l00147"></a><span class="lineno">  147</span>&#160;            }</div><div class="line"><a name="l00148"></a><span class="lineno">  148</span>&#160;        }</div><div class="line"><a name="l00149"></a><span class="lineno">  149</span>&#160;    }</div><div class="line"><a name="l00150"></a><span class="lineno">  150</span>&#160;<span class="preprocessor">#if defined HAS_BIAS</span></div><div class="line"><a name="l00151"></a><span class="lineno">  151</span>&#160;    *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = (<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>)(1);</div><div class="line"><a name="l00152"></a><span class="lineno">  152</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00153"></a><span class="lineno">  153</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00079">helpers.h:79</a></div></div>
 448 <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>
 449 <div class="ttc" id="helpers_8h_xhtml_a4334a4a76f8e9628c0fb9e1acf616e2a"><div class="ttname"><a href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00076">helpers.h:76</a></div></div>
 450 <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#l00104">helpers.h:104</a></div></div>
 451 <div class="ttc" id="helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00213">helpers.h:213</a></div></div>
 452 <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#l00095">helpers.h:95</a></div></div>
 453 <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#l00097">helpers.h:97</a></div></div>
 454 <div class="ttc" id="struct_image_xhtml_a4f0b90c9ecd6e57ceb3f37332fefe8f1"><div class="ttname"><a href="struct_image.xhtml#a4f0b90c9ecd6e57ceb3f37332fefe8f1">Image::stride_y</a></div><div class="ttdeci">int stride_y</div><div class="ttdoc">Stride of the image in Y dimension (in bytes) </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00100">helpers.h:100</a></div></div>
 455 <div class="ttc" id="struct_image_xhtml_ae01febbfd0689ef709f3ff6fdd2abc7e"><div class="ttname"><a href="struct_image.xhtml#ae01febbfd0689ef709f3ff6fdd2abc7e">Image::stride_x</a></div><div class="ttdeci">int stride_x</div><div class="ttdoc">Stride of the image in X dimension (in bytes) </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
 456 </div><!-- fragment -->
 457 </div>
 458 </div>
 459 <a class="anchor" id="a6f31f2c8d606e1b57d9bdfacd416024c"></a>
 460 <div class="memitem">
 461 <div class="memproto">
 462       <table class="memname">
 463         <tr>
 464           <td class="memname">__kernel void im2col_reduced </td>
 465           <td>(</td>
 466           <td class="paramtype">__global uchar *&#160;</td>
 467           <td class="paramname"><em>src_ptr</em>, </td>
 468         </tr>
 469         <tr>
 470           <td class="paramkey"></td>
 471           <td></td>
 472           <td class="paramtype">uint&#160;</td>
 473           <td class="paramname"><em>src_stride_x</em>, </td>
 474         </tr>
 475         <tr>
 476           <td class="paramkey"></td>
 477           <td></td>
 478           <td class="paramtype">uint&#160;</td>
 479           <td class="paramname"><em>src_step_x</em>, </td>
 480         </tr>
 481         <tr>
 482           <td class="paramkey"></td>
 483           <td></td>
 484           <td class="paramtype">uint&#160;</td>
 485           <td class="paramname"><em>src_stride_y</em>, </td>
 486         </tr>
 487         <tr>
 488           <td class="paramkey"></td>
 489           <td></td>
 490           <td class="paramtype">uint&#160;</td>
 491           <td class="paramname"><em>src_step_y</em>, </td>
 492         </tr>
 493         <tr>
 494           <td class="paramkey"></td>
 495           <td></td>
 496           <td class="paramtype">uint&#160;</td>
 497           <td class="paramname"><em>src_stride_z</em>, </td>
 498         </tr>
 499         <tr>
 500           <td class="paramkey"></td>
 501           <td></td>
 502           <td class="paramtype">uint&#160;</td>
 503           <td class="paramname"><em>src_step_z</em>, </td>
 504         </tr>
 505         <tr>
 506           <td class="paramkey"></td>
 507           <td></td>
 508           <td class="paramtype">uint&#160;</td>
 509           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 510         </tr>
 511         <tr>
 512           <td class="paramkey"></td>
 513           <td></td>
 514           <td class="paramtype">__global uchar *&#160;</td>
 515           <td class="paramname"><em>dst_ptr</em>, </td>
 516         </tr>
 517         <tr>
 518           <td class="paramkey"></td>
 519           <td></td>
 520           <td class="paramtype">uint&#160;</td>
 521           <td class="paramname"><em>dst_stride_x</em>, </td>
 522         </tr>
 523         <tr>
 524           <td class="paramkey"></td>
 525           <td></td>
 526           <td class="paramtype">uint&#160;</td>
 527           <td class="paramname"><em>dst_step_x</em>, </td>
 528         </tr>
 529         <tr>
 530           <td class="paramkey"></td>
 531           <td></td>
 532           <td class="paramtype">uint&#160;</td>
 533           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
 534         </tr>
 535         <tr>
 536           <td class="paramkey"></td>
 537           <td></td>
 538           <td class="paramtype">uint&#160;</td>
 539           <td class="paramname"><em>width</em>, </td>
 540         </tr>
 541         <tr>
 542           <td class="paramkey"></td>
 543           <td></td>
 544           <td class="paramtype">uint&#160;</td>
 545           <td class="paramname"><em>height</em>&#160;</td>
 546         </tr>
 547         <tr>
 548           <td></td>
 549           <td>)</td>
 550           <td></td><td></td>
 551         </tr>
 552       </table>
 553 </div><div class="memdoc">
 554
 555 <p>This kernel reshapes the tensor's low three dimensions to single row for GEMM operation. </p>
 556 <dl class="section note"><dt>Note</dt><dd>Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float </dd>
 557 <dd>
 558 In case biases will be added in late stage, -DHAS_BIAS has to be passed to append the final matrix with 1 in each row.</dd></dl>
 559 <dl class="params"><dt>Parameters</dt><dd>
 560   <table class="params">
 561     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: F16, F32 </td></tr>
 562     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source tensor in X dimension (in bytes) </td></tr>
 563     <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>
 564     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source tensor in Y dimension (in bytes) </td></tr>
 565     <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>
 566     <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>
 567     <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>
 568     <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 tensor </td></tr>
 569     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Same as input. </td></tr>
 570     <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>
 571     <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>
 572     <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>
 573     <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>The width of the input tensor </td></tr>
 574     <tr><td class="paramdir">[in]</td><td class="paramname">height</td><td>The height of the input tensor </td></tr>
 575   </table>
 576   </dd>
 577 </dl>
 578
 579 <p>Definition at line <a class="el" href="convolution__layer_8cl_source.xhtml#l00208">208</a> of file <a class="el" href="convolution__layer_8cl_source.xhtml">convolution_layer.cl</a>.</p>
 580
 581 <p>References <a class="el" href="helpers_8h_source.xhtml#l00079">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, and <a class="el" href="helpers_8h_source.xhtml#l00106">Tensor3D::ptr</a>.</p>
 582 <div class="fragment"><div class="line"><a name="l00212"></a><span class="lineno">  212</span>&#160;{</div><div class="line"><a name="l00213"></a><span class="lineno">  213</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> src = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(src);</div><div class="line"><a name="l00214"></a><span class="lineno">  214</span>&#160;</div><div class="line"><a name="l00215"></a><span class="lineno">  215</span>&#160;    <span class="keyword">const</span> uint image_size = width * height;</div><div class="line"><a name="l00216"></a><span class="lineno">  216</span>&#160;</div><div class="line"><a name="l00217"></a><span class="lineno">  217</span>&#160;    __global uchar *tmp_out_ptr = dst_ptr + dst_offset_first_element_in_bytes + (get_global_id(0) + get_global_id(1) * width + get_global_id(2) * image_size) * dst_stride_x;</div><div class="line"><a name="l00218"></a><span class="lineno">  218</span>&#160;</div><div class="line"><a name="l00219"></a><span class="lineno">  219</span>&#160;    *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)src.ptr);</div><div class="line"><a name="l00220"></a><span class="lineno">  220</span>&#160;</div><div class="line"><a name="l00221"></a><span class="lineno">  221</span>&#160;<span class="preprocessor">#if defined HAS_BIAS</span></div><div class="line"><a name="l00222"></a><span class="lineno">  222</span>&#160;    <span class="comment">// If it is the last thread in the 3 dimensional workgroup</span></div><div class="line"><a name="l00223"></a><span class="lineno">  223</span>&#160;    <span class="keywordflow">if</span>(get_global_id(0) == (get_global_size(0) - 1) &amp;&amp; get_global_id(1) == (get_global_size(1) - 1) &amp;&amp; get_global_id(2) == (get_global_size(2) - 1))</div><div class="line"><a name="l00224"></a><span class="lineno">  224</span>&#160;    {</div><div class="line"><a name="l00225"></a><span class="lineno">  225</span>&#160;        tmp_out_ptr += dst_stride_x;</div><div class="line"><a name="l00226"></a><span class="lineno">  226</span>&#160;        *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = (<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>)1;</div><div class="line"><a name="l00227"></a><span class="lineno">  227</span>&#160;    }</div><div class="line"><a name="l00228"></a><span class="lineno">  228</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00229"></a><span class="lineno">  229</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00079">helpers.h:79</a></div></div>
 583 <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>
 584 <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#l00104">helpers.h:104</a></div></div>
 585 </div><!-- fragment -->
 586 </div>
 587 </div>
 588 <a class="anchor" id="a229c10cbfba415bce4cc33d4530bfbd0"></a>
 589 <div class="memitem">
 590 <div class="memproto">
 591       <table class="memname">
 592         <tr>
 593           <td class="memname">__kernel void reshape_to_columns </td>
 594           <td>(</td>
 595           <td class="paramtype">__global uchar *&#160;</td>
 596           <td class="paramname"><em>src_ptr</em>, </td>
 597         </tr>
 598         <tr>
 599           <td class="paramkey"></td>
 600           <td></td>
 601           <td class="paramtype">uint&#160;</td>
 602           <td class="paramname"><em>src_stride_x</em>, </td>
 603         </tr>
 604         <tr>
 605           <td class="paramkey"></td>
 606           <td></td>
 607           <td class="paramtype">uint&#160;</td>
 608           <td class="paramname"><em>src_step_x</em>, </td>
 609         </tr>
 610         <tr>
 611           <td class="paramkey"></td>
 612           <td></td>
 613           <td class="paramtype">uint&#160;</td>
 614           <td class="paramname"><em>src_stride_y</em>, </td>
 615         </tr>
 616         <tr>
 617           <td class="paramkey"></td>
 618           <td></td>
 619           <td class="paramtype">uint&#160;</td>
 620           <td class="paramname"><em>src_step_y</em>, </td>
 621         </tr>
 622         <tr>
 623           <td class="paramkey"></td>
 624           <td></td>
 625           <td class="paramtype">uint&#160;</td>
 626           <td class="paramname"><em>src_stride_z</em>, </td>
 627         </tr>
 628         <tr>
 629           <td class="paramkey"></td>
 630           <td></td>
 631           <td class="paramtype">uint&#160;</td>
 632           <td class="paramname"><em>src_step_z</em>, </td>
 633         </tr>
 634         <tr>
 635           <td class="paramkey"></td>
 636           <td></td>
 637           <td class="paramtype">uint&#160;</td>
 638           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 639         </tr>
 640         <tr>
 641           <td class="paramkey"></td>
 642           <td></td>
 643           <td class="paramtype">__global uchar *&#160;</td>
 644           <td class="paramname"><em>dst_ptr</em>, </td>
 645         </tr>
 646         <tr>
 647           <td class="paramkey"></td>
 648           <td></td>
 649           <td class="paramtype">uint&#160;</td>
 650           <td class="paramname"><em>dst_stride_x</em>, </td>
 651         </tr>
 652         <tr>
 653           <td class="paramkey"></td>
 654           <td></td>
 655           <td class="paramtype">uint&#160;</td>
 656           <td class="paramname"><em>dst_step_x</em>, </td>
 657         </tr>
 658         <tr>
 659           <td class="paramkey"></td>
 660           <td></td>
 661           <td class="paramtype">uint&#160;</td>
 662           <td class="paramname"><em>dst_stride_y</em>, </td>
 663         </tr>
 664         <tr>
 665           <td class="paramkey"></td>
 666           <td></td>
 667           <td class="paramtype">uint&#160;</td>
 668           <td class="paramname"><em>dst_step_y</em>, </td>
 669         </tr>
 670         <tr>
 671           <td class="paramkey"></td>
 672           <td></td>
 673           <td class="paramtype">uint&#160;</td>
 674           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
 675         </tr>
 676         <tr>
 677           <td class="paramkey"></td>
 678           <td></td>
 679           <td class="paramtype">__global uchar *&#160;</td>
 680           <td class="paramname"><em>bias_ptr</em>, </td>
 681         </tr>
 682         <tr>
 683           <td class="paramkey"></td>
 684           <td></td>
 685           <td class="paramtype">uint&#160;</td>
 686           <td class="paramname"><em>bias_stride_x</em>, </td>
 687         </tr>
 688         <tr>
 689           <td class="paramkey"></td>
 690           <td></td>
 691           <td class="paramtype">uint&#160;</td>
 692           <td class="paramname"><em>bias_step_x</em>, </td>
 693         </tr>
 694         <tr>
 695           <td class="paramkey"></td>
 696           <td></td>
 697           <td class="paramtype">uint&#160;</td>
 698           <td class="paramname"><em>bias_stride_y</em>, </td>
 699         </tr>
 700         <tr>
 701           <td class="paramkey"></td>
 702           <td></td>
 703           <td class="paramtype">uint&#160;</td>
 704           <td class="paramname"><em>bias_step_y</em>, </td>
 705         </tr>
 706         <tr>
 707           <td class="paramkey"></td>
 708           <td></td>
 709           <td class="paramtype">uint&#160;</td>
 710           <td class="paramname"><em>bias_offset_first_element_in_bytes</em>, </td>
 711         </tr>
 712         <tr>
 713           <td class="paramkey"></td>
 714           <td></td>
 715           <td class="paramtype">uint&#160;</td>
 716           <td class="paramname"><em>width</em>, </td>
 717         </tr>
 718         <tr>
 719           <td class="paramkey"></td>
 720           <td></td>
 721           <td class="paramtype">uint&#160;</td>
 722           <td class="paramname"><em>height</em>&#160;</td>
 723         </tr>
 724         <tr>
 725           <td></td>
 726           <td>)</td>
 727           <td></td><td></td>
 728         </tr>
 729       </table>
 730 </div><div class="memdoc">
 731
 732 <p>This kernel reshapes the tensor's low three dimensions to single column. </p>
 733 <dl class="section note"><dt>Note</dt><dd>Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short</dd></dl>
 734 <dl class="params"><dt>Parameters</dt><dd>
 735   <table class="params">
 736     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: F16, F32 </td></tr>
 737     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source tensor in X dimension (in bytes) </td></tr>
 738     <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>
 739     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source tensor in Y dimension (in bytes) </td></tr>
 740     <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>
 741     <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>
 742     <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>
 743     <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 tensor </td></tr>
 744     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Same as input </td></tr>
 745     <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>
 746     <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>
 747     <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>
 748     <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>
 749     <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>
 750     <tr><td class="paramdir">[in]</td><td class="paramname">bias_ptr</td><td>Pointer to the bias tensor. Same as input </td></tr>
 751     <tr><td class="paramdir">[in]</td><td class="paramname">bias_stride_x</td><td>Stride of the bias tensor in X dimension (in bytes) </td></tr>
 752     <tr><td class="paramdir">[in]</td><td class="paramname">bias_step_x</td><td>bias_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 753     <tr><td class="paramdir">[in]</td><td class="paramname">bias_stride_y</td><td>Stride of the bias tensor in Y dimension (in bytes) </td></tr>
 754     <tr><td class="paramdir">[in]</td><td class="paramname">bias_step_y</td><td>bias_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 755     <tr><td class="paramdir">[in]</td><td class="paramname">bias_offset_first_element_in_bytes</td><td>The offset of the first element in the source tensor </td></tr>
 756     <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>The width of the input tensor </td></tr>
 757     <tr><td class="paramdir">[in]</td><td class="paramname">height</td><td>The height of the input tensor </td></tr>
 758   </table>
 759   </dd>
 760 </dl>
 761
 762 <p>Definition at line <a class="el" href="convolution__layer_8cl_source.xhtml#l00053">53</a> of file <a class="el" href="convolution__layer_8cl_source.xhtml">convolution_layer.cl</a>.</p>
 763
 764 <p>References <a class="el" href="helpers_8h_source.xhtml#l00079">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, and <a class="el" href="helpers_8h_source.xhtml#l00106">Tensor3D::ptr</a>.</p>
 765 <div class="fragment"><div class="line"><a name="l00060"></a><span class="lineno">   60</span>&#160;{</div><div class="line"><a name="l00061"></a><span class="lineno">   61</span>&#160;    <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> src = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(src);</div><div class="line"><a name="l00062"></a><span class="lineno">   62</span>&#160;</div><div class="line"><a name="l00063"></a><span class="lineno">   63</span>&#160;    __global uchar *tmp_out_ptr = dst_ptr + dst_offset_first_element_in_bytes + get_global_id(0) * dst_stride_y</div><div class="line"><a name="l00064"></a><span class="lineno">   64</span>&#160;                                  + get_global_id(1) * width * dst_stride_y + get_global_id(2) * width * height * dst_stride_y;</div><div class="line"><a name="l00065"></a><span class="lineno">   65</span>&#160;</div><div class="line"><a name="l00066"></a><span class="lineno">   66</span>&#160;    *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)src.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00067"></a><span class="lineno">   67</span>&#160;</div><div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;<span class="preprocessor">#if defined HAS_BIAS</span></div><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;    <span class="comment">// If it is the last thread in the 3 dimensional workgroup</span></div><div class="line"><a name="l00070"></a><span class="lineno">   70</span>&#160;    <span class="keywordflow">if</span>(get_global_id(0) == (get_global_size(0) - 1) &amp;&amp; get_global_id(1) == (get_global_size(1) - 1) &amp;&amp; get_global_id(2) == (get_global_size(2) - 1))</div><div class="line"><a name="l00071"></a><span class="lineno">   71</span>&#160;    {</div><div class="line"><a name="l00072"></a><span class="lineno">   72</span>&#160;        tmp_out_ptr += dst_stride_y;</div><div class="line"><a name="l00073"></a><span class="lineno">   73</span>&#160;        *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)tmp_out_ptr) = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)(bias_ptr + bias_offset_first_element_in_bytes));</div><div class="line"><a name="l00074"></a><span class="lineno">   74</span>&#160;    }</div><div class="line"><a name="l00075"></a><span class="lineno">   75</span>&#160;</div><div class="line"><a name="l00076"></a><span class="lineno">   76</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00079">helpers.h:79</a></div></div>
 766 <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>
 767 <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#l00104">helpers.h:104</a></div></div>
 768 <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#l00106">helpers.h:106</a></div></div>
 769 </div><!-- fragment -->
 770 </div>
 771 </div>
 772 </div><!-- contents -->
 773 </div><!-- doc-content -->
 774 <!-- start footer part -->
 775 <div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
 776   <ul>
 777     <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="convolution__layer_8cl.xhtml">convolution_layer.cl</a></li>
 778     <li class="footer">Generated on Wed Apr 12 2017 14:26:05 for ARM Compute Library by
 779     <a href="http://www.doxygen.org/index.html">
 780     <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
 781   </ul>
 782 </div>
 783 </body>
 784 </html>