documentation/gemm_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/gemm.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('gemm_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">gemm.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="gemm_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:a21d7061f36aec78824e368f2ef1cafc1"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a21d7061f36aec78824e368f2ef1cafc1">gemm_transpose1x4_f32</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_offset_first_element_in_bytes)</td></tr>
 128 <tr class="memdesc:a21d7061f36aec78824e368f2ef1cafc1"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel computes the "vector" 1x4 transposition of input matrix.  <a href="#a21d7061f36aec78824e368f2ef1cafc1">More...</a><br /></td></tr>
 129 <tr class="separator:a21d7061f36aec78824e368f2ef1cafc1"><td class="memSeparator" colspan="2">&#160;</td></tr>
 130 <tr class="memitem:a11d87ecaf1388b4f57404e062265b374"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a11d87ecaf1388b4f57404e062265b374">gemm_transpose1x8_f16</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_offset_first_element_in_bytes)</td></tr>
 131 <tr class="memdesc:a11d87ecaf1388b4f57404e062265b374"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel computes the "vector" 1x8 transposition of input matrix.  <a href="#a11d87ecaf1388b4f57404e062265b374">More...</a><br /></td></tr>
 132 <tr class="separator:a11d87ecaf1388b4f57404e062265b374"><td class="memSeparator" colspan="2">&#160;</td></tr>
 133 <tr class="memitem:a675879eb72964ddbd5f71155da3b0cc3"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a675879eb72964ddbd5f71155da3b0cc3">gemm_transpose1x16_u8</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_offset_first_element_in_bytes)</td></tr>
 134 <tr class="memdesc:a675879eb72964ddbd5f71155da3b0cc3"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel computes the "vector" 1x16 transposition of input matrix.  <a href="#a675879eb72964ddbd5f71155da3b0cc3">More...</a><br /></td></tr>
 135 <tr class="separator:a675879eb72964ddbd5f71155da3b0cc3"><td class="memSeparator" colspan="2">&#160;</td></tr>
 136 <tr class="memitem:acd423b6f992354e7c00137b20d687281"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#acd423b6f992354e7c00137b20d687281">gemm_interleave4x4_f32</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_offset_first_element_in_bytes)</td></tr>
 137 <tr class="memdesc:acd423b6f992354e7c00137b20d687281"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values.  <a href="#acd423b6f992354e7c00137b20d687281">More...</a><br /></td></tr>
 138 <tr class="separator:acd423b6f992354e7c00137b20d687281"><td class="memSeparator" colspan="2">&#160;</td></tr>
 139 <tr class="memitem:ae333c12d780666b2591f7c8e9faaf5a9"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#ae333c12d780666b2591f7c8e9faaf5a9">gemm_interleave4x4_f16</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_offset_first_element_in_bytes)</td></tr>
 140 <tr class="memdesc:ae333c12d780666b2591f7c8e9faaf5a9"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values.  <a href="#ae333c12d780666b2591f7c8e9faaf5a9">More...</a><br /></td></tr>
 141 <tr class="separator:ae333c12d780666b2591f7c8e9faaf5a9"><td class="memSeparator" colspan="2">&#160;</td></tr>
 142 <tr class="memitem:a830ba1cc0ad3c8cffa4a14424b2d0411"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a830ba1cc0ad3c8cffa4a14424b2d0411">gemm_interleave4x4_u8</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_offset_first_element_in_bytes)</td></tr>
 143 <tr class="memdesc:a830ba1cc0ad3c8cffa4a14424b2d0411"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values.  <a href="#a830ba1cc0ad3c8cffa4a14424b2d0411">More...</a><br /></td></tr>
 144 <tr class="separator:a830ba1cc0ad3c8cffa4a14424b2d0411"><td class="memSeparator" colspan="2">&#160;</td></tr>
 145 <tr class="memitem:a7c619984bc9c31f5f6c19ecc88b56581"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a7c619984bc9c31f5f6c19ecc88b56581">gemm_accumulate_biases_f32</a> (__global uchar *accum_ptr, uint accum_stride_x, uint accum_step_x, uint accum_stride_y, uint accum_step_y, uint accum_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>
 146 <tr class="memdesc:a7c619984bc9c31f5f6c19ecc88b56581"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel accumulates each row with the biases vector.  <a href="#a7c619984bc9c31f5f6c19ecc88b56581">More...</a><br /></td></tr>
 147 <tr class="separator:a7c619984bc9c31f5f6c19ecc88b56581"><td class="memSeparator" colspan="2">&#160;</td></tr>
 148 <tr class="memitem:a76295c333dfa3c9bb20975579dce87e8"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a76295c333dfa3c9bb20975579dce87e8">gemm_accumulate_biases_f16</a> (__global uchar *accum_ptr, uint accum_stride_x, uint accum_step_x, uint accum_stride_y, uint accum_step_y, uint accum_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>
 149 <tr class="memdesc:a76295c333dfa3c9bb20975579dce87e8"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel accumulates each row with the biases vector.  <a href="#a76295c333dfa3c9bb20975579dce87e8">More...</a><br /></td></tr>
 150 <tr class="separator:a76295c333dfa3c9bb20975579dce87e8"><td class="memSeparator" colspan="2">&#160;</td></tr>
 151 </table>
 152 <h2 class="groupheader">Function Documentation</h2>
 153 <a class="anchor" id="a76295c333dfa3c9bb20975579dce87e8"></a>
 154 <div class="memitem">
 155 <div class="memproto">
 156       <table class="memname">
 157         <tr>
 158           <td class="memname">__kernel void gemm_accumulate_biases_f16 </td>
 159           <td>(</td>
 160           <td class="paramtype">__global uchar *&#160;</td>
 161           <td class="paramname"><em>accum_ptr</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>accum_stride_x</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>accum_step_x</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>accum_stride_y</em>, </td>
 180         </tr>
 181         <tr>
 182           <td class="paramkey"></td>
 183           <td></td>
 184           <td class="paramtype">uint&#160;</td>
 185           <td class="paramname"><em>accum_step_y</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>accum_offset_first_element_in_bytes</em>, </td>
 192         </tr>
 193         <tr>
 194           <td class="paramkey"></td>
 195           <td></td>
 196           <td class="paramtype">__global uchar *&#160;</td>
 197           <td class="paramname"><em>biases_ptr</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>biases_stride_x</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>biases_step_x</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>biases_offset_first_element_in_bytes</em>&#160;</td>
 216         </tr>
 217         <tr>
 218           <td></td>
 219           <td>)</td>
 220           <td></td><td></td>
 221         </tr>
 222       </table>
 223 </div><div class="memdoc">
 224
 225 <p>This kernel accumulates each row with the biases vector. </p>
 226 <dl class="params"><dt>Parameters</dt><dd>
 227   <table class="params">
 228     <tr><td class="paramdir">[in,out]</td><td class="paramname">accum_ptr</td><td>Pointer to the accumulate tensor. Supported data type: F16 </td></tr>
 229     <tr><td class="paramdir">[in]</td><td class="paramname">accum_stride_x</td><td>Stride of the accumulate tensor in X dimension (in bytes) </td></tr>
 230     <tr><td class="paramdir">[in]</td><td class="paramname">accum_step_x</td><td>accum_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 231     <tr><td class="paramdir">[in]</td><td class="paramname">accum_stride_y</td><td>Stride of the accumlulate tensor in Y dimension (in bytes) </td></tr>
 232     <tr><td class="paramdir">[in]</td><td class="paramname">accum_step_y</td><td>src_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 233     <tr><td class="paramdir">[in]</td><td class="paramname">accum_offset_first_element_in_bytes</td><td>The offset of the first element in the accumulate tensor </td></tr>
 234     <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>Pointer to the biases vector. Same as input. </td></tr>
 235     <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>Stride of the destination tensor in X dimension (in bytes) </td></tr>
 236     <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>dst_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 237     <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>The offset of the first element in the destination tensor </td></tr>
 238   </table>
 239   </dd>
 240 </dl>
 241
 242 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00290">290</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
 243
 244 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00067">CONVERT_TO_VECTOR_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00049">IMAGE_DECLARATION</a>, <a class="el" href="helpers_8h_source.xhtml#l00201">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00089">Vector::ptr</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
 245 <div class="fragment"><div class="line"><a name="l00293"></a><span class="lineno">  293</span>&#160;{</div><div class="line"><a name="l00294"></a><span class="lineno">  294</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>  accum  = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(accum);</div><div class="line"><a name="l00295"></a><span class="lineno">  295</span>&#160;    <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a>(biases);</div><div class="line"><a name="l00296"></a><span class="lineno">  296</span>&#160;</div><div class="line"><a name="l00297"></a><span class="lineno">  297</span>&#160;    half8 accum_value  = vload8(0, (__global half *)accum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00298"></a><span class="lineno">  298</span>&#160;    half8 biases_value = vload8(0, (__global half *)biases.<a class="code" href="struct_vector.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00299"></a><span class="lineno">  299</span>&#160;    accum_value        = biases_value + accum_value;</div><div class="line"><a name="l00300"></a><span class="lineno">  300</span>&#160;</div><div class="line"><a name="l00301"></a><span class="lineno">  301</span>&#160;    <span class="comment">// Store result in the accummulate buffer</span></div><div class="line"><a name="l00302"></a><span class="lineno">  302</span>&#160;    vstore8(accum_value, 0, (__global half *)accum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00303"></a><span class="lineno">  303</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#l00087">helpers.h:87</a></div></div>
 246 <div class="ttc" id="helpers_8h_xhtml_a527bfdf5eeb306f1cf01c4a8e29f38e0"><div class="ttname"><a href="helpers_8h.xhtml#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00067">helpers.h:67</a></div></div>
 247 <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>
 248 <div class="ttc" id="struct_vector_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_vector.xhtml#acf52c23cbd7424606c10a606524e3e32">Vector::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#l00089">helpers.h:89</a></div></div>
 249 <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>
 250 <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>
 251 </div><!-- fragment -->
 252 </div>
 253 </div>
 254 <a class="anchor" id="a7c619984bc9c31f5f6c19ecc88b56581"></a>
 255 <div class="memitem">
 256 <div class="memproto">
 257       <table class="memname">
 258         <tr>
 259           <td class="memname">__kernel void gemm_accumulate_biases_f32 </td>
 260           <td>(</td>
 261           <td class="paramtype">__global uchar *&#160;</td>
 262           <td class="paramname"><em>accum_ptr</em>, </td>
 263         </tr>
 264         <tr>
 265           <td class="paramkey"></td>
 266           <td></td>
 267           <td class="paramtype">uint&#160;</td>
 268           <td class="paramname"><em>accum_stride_x</em>, </td>
 269         </tr>
 270         <tr>
 271           <td class="paramkey"></td>
 272           <td></td>
 273           <td class="paramtype">uint&#160;</td>
 274           <td class="paramname"><em>accum_step_x</em>, </td>
 275         </tr>
 276         <tr>
 277           <td class="paramkey"></td>
 278           <td></td>
 279           <td class="paramtype">uint&#160;</td>
 280           <td class="paramname"><em>accum_stride_y</em>, </td>
 281         </tr>
 282         <tr>
 283           <td class="paramkey"></td>
 284           <td></td>
 285           <td class="paramtype">uint&#160;</td>
 286           <td class="paramname"><em>accum_step_y</em>, </td>
 287         </tr>
 288         <tr>
 289           <td class="paramkey"></td>
 290           <td></td>
 291           <td class="paramtype">uint&#160;</td>
 292           <td class="paramname"><em>accum_offset_first_element_in_bytes</em>, </td>
 293         </tr>
 294         <tr>
 295           <td class="paramkey"></td>
 296           <td></td>
 297           <td class="paramtype">__global uchar *&#160;</td>
 298           <td class="paramname"><em>biases_ptr</em>, </td>
 299         </tr>
 300         <tr>
 301           <td class="paramkey"></td>
 302           <td></td>
 303           <td class="paramtype">uint&#160;</td>
 304           <td class="paramname"><em>biases_stride_x</em>, </td>
 305         </tr>
 306         <tr>
 307           <td class="paramkey"></td>
 308           <td></td>
 309           <td class="paramtype">uint&#160;</td>
 310           <td class="paramname"><em>biases_step_x</em>, </td>
 311         </tr>
 312         <tr>
 313           <td class="paramkey"></td>
 314           <td></td>
 315           <td class="paramtype">uint&#160;</td>
 316           <td class="paramname"><em>biases_offset_first_element_in_bytes</em>&#160;</td>
 317         </tr>
 318         <tr>
 319           <td></td>
 320           <td>)</td>
 321           <td></td><td></td>
 322         </tr>
 323       </table>
 324 </div><div class="memdoc">
 325
 326 <p>This kernel accumulates each row with the biases vector. </p>
 327 <dl class="params"><dt>Parameters</dt><dd>
 328   <table class="params">
 329     <tr><td class="paramdir">[in,out]</td><td class="paramname">accum_ptr</td><td>Pointer to the accumulate tensor. Supported data type: F32 </td></tr>
 330     <tr><td class="paramdir">[in]</td><td class="paramname">accum_stride_x</td><td>Stride of the accmulate tensor in X dimension (in bytes) </td></tr>
 331     <tr><td class="paramdir">[in]</td><td class="paramname">accum_step_x</td><td>accum_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 332     <tr><td class="paramdir">[in]</td><td class="paramname">accum_stride_y</td><td>Stride of the accumlulate tensor in Y dimension (in bytes) </td></tr>
 333     <tr><td class="paramdir">[in]</td><td class="paramname">accum_step_y</td><td>src_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 334     <tr><td class="paramdir">[in]</td><td class="paramname">accum_offset_first_element_in_bytes</td><td>The offset of the first element in the accumulate tensor </td></tr>
 335     <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>Pointer to the biases vector. Same as input. </td></tr>
 336     <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>Stride of the destination tensor in X dimension (in bytes) </td></tr>
 337     <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>dst_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 338     <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>The offset of the first element in the destination tensor </td></tr>
 339   </table>
 340   </dd>
 341 </dl>
 342
 343 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00262">262</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
 344
 345 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00067">CONVERT_TO_VECTOR_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00089">Vector::ptr</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
 346 <div class="fragment"><div class="line"><a name="l00265"></a><span class="lineno">  265</span>&#160;{</div><div class="line"><a name="l00266"></a><span class="lineno">  266</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a>  accum  = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(accum);</div><div class="line"><a name="l00267"></a><span class="lineno">  267</span>&#160;    <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a>(biases);</div><div class="line"><a name="l00268"></a><span class="lineno">  268</span>&#160;</div><div class="line"><a name="l00269"></a><span class="lineno">  269</span>&#160;    float4 accum_value  = vload4(0, (__global <span class="keywordtype">float</span> *)accum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00270"></a><span class="lineno">  270</span>&#160;    float4 biases_value = vload4(0, (__global <span class="keywordtype">float</span> *)biases.<a class="code" href="struct_vector.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00271"></a><span class="lineno">  271</span>&#160;    accum_value         = biases_value + accum_value;</div><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;    <span class="comment">// Store result in the accummulate buffer</span></div><div class="line"><a name="l00274"></a><span class="lineno">  274</span>&#160;    vstore4(accum_value, 0, (__global <span class="keywordtype">float</span> *)accum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00275"></a><span class="lineno">  275</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#l00087">helpers.h:87</a></div></div>
 347 <div class="ttc" id="helpers_8h_xhtml_a527bfdf5eeb306f1cf01c4a8e29f38e0"><div class="ttname"><a href="helpers_8h.xhtml#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00067">helpers.h:67</a></div></div>
 348 <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>
 349 <div class="ttc" id="struct_vector_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_vector.xhtml#acf52c23cbd7424606c10a606524e3e32">Vector::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#l00089">helpers.h:89</a></div></div>
 350 <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>
 351 <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>
 352 </div><!-- fragment -->
 353 </div>
 354 </div>
 355 <a class="anchor" id="ae333c12d780666b2591f7c8e9faaf5a9"></a>
 356 <div class="memitem">
 357 <div class="memproto">
 358       <table class="memname">
 359         <tr>
 360           <td class="memname">__kernel void gemm_interleave4x4_f16 </td>
 361           <td>(</td>
 362           <td class="paramtype">__global uchar *&#160;</td>
 363           <td class="paramname"><em>src_ptr</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>src_stride_x</em>, </td>
 370         </tr>
 371         <tr>
 372           <td class="paramkey"></td>
 373           <td></td>
 374           <td class="paramtype">uint&#160;</td>
 375           <td class="paramname"><em>src_step_x</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>src_stride_y</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>src_step_y</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>src_offset_first_element_in_bytes</em>, </td>
 394         </tr>
 395         <tr>
 396           <td class="paramkey"></td>
 397           <td></td>
 398           <td class="paramtype">__global uchar *&#160;</td>
 399           <td class="paramname"><em>dst_ptr</em>, </td>
 400         </tr>
 401         <tr>
 402           <td class="paramkey"></td>
 403           <td></td>
 404           <td class="paramtype">uint&#160;</td>
 405           <td class="paramname"><em>dst_stride_x</em>, </td>
 406         </tr>
 407         <tr>
 408           <td class="paramkey"></td>
 409           <td></td>
 410           <td class="paramtype">uint&#160;</td>
 411           <td class="paramname"><em>dst_step_x</em>, </td>
 412         </tr>
 413         <tr>
 414           <td class="paramkey"></td>
 415           <td></td>
 416           <td class="paramtype">uint&#160;</td>
 417           <td class="paramname"><em>dst_stride_y</em>, </td>
 418         </tr>
 419         <tr>
 420           <td class="paramkey"></td>
 421           <td></td>
 422           <td class="paramtype">uint&#160;</td>
 423           <td class="paramname"><em>dst_step_y</em>, </td>
 424         </tr>
 425         <tr>
 426           <td class="paramkey"></td>
 427           <td></td>
 428           <td class="paramtype">uint&#160;</td>
 429           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
 430         </tr>
 431         <tr>
 432           <td></td>
 433           <td>)</td>
 434           <td></td><td></td>
 435         </tr>
 436       </table>
 437 </div><div class="memdoc">
 438
 439 <p>This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. </p>
 440 <dl class="params"><dt>Parameters</dt><dd>
 441   <table class="params">
 442     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: F16 </td></tr>
 443     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
 444     <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>
 445     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
 446     <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>
 447     <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 matrix </td></tr>
 448     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: F16 </td></tr>
 449     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
 450     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 451     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
 452     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 453     <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 matrix </td></tr>
 454   </table>
 455   </dd>
 456 </dl>
 457
 458 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00178">178</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
 459
 460 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00201">offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
 461 <div class="fragment"><div class="line"><a name="l00180"></a><span class="lineno">  180</span>&#160;{</div><div class="line"><a name="l00181"></a><span class="lineno">  181</span>&#160;    <span class="comment">/* Compute source and destination addresses */</span></div><div class="line"><a name="l00182"></a><span class="lineno">  182</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="l00183"></a><span class="lineno">  183</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> dst = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00184"></a><span class="lineno">  184</span>&#160;</div><div class="line"><a name="l00185"></a><span class="lineno">  185</span>&#160;    <span class="comment">/* Load values from Matrix A */</span></div><div class="line"><a name="l00186"></a><span class="lineno">  186</span>&#160;    half8 a0 = vload8(0, (__global half *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0)));</div><div class="line"><a name="l00187"></a><span class="lineno">  187</span>&#160;    half8 a1 = vload8(0, (__global half *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 1)));</div><div class="line"><a name="l00188"></a><span class="lineno">  188</span>&#160;    half8 a2 = vload8(0, (__global half *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 2)));</div><div class="line"><a name="l00189"></a><span class="lineno">  189</span>&#160;    half8 a3 = vload8(0, (__global half *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 3)));</div><div class="line"><a name="l00190"></a><span class="lineno">  190</span>&#160;</div><div class="line"><a name="l00191"></a><span class="lineno">  191</span>&#160;    half8 val0 = (half8)((half4)(a0.s0, a1.s0, a2.s0, a3.s0), (half4)(a0.s1, a1.s1, a2.s1, a3.s1));</div><div class="line"><a name="l00192"></a><span class="lineno">  192</span>&#160;    vstore8(val0, 0, ((__global half *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 0);</div><div class="line"><a name="l00193"></a><span class="lineno">  193</span>&#160;</div><div class="line"><a name="l00194"></a><span class="lineno">  194</span>&#160;    val0 = (half8)((half4)(a0.s2, a1.s2, a2.s2, a3.s2), (half4)(a0.s3, a1.s3, a2.s3, a3.s3));</div><div class="line"><a name="l00195"></a><span class="lineno">  195</span>&#160;    vstore8(val0, 0, ((__global half *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 8);</div><div class="line"><a name="l00196"></a><span class="lineno">  196</span>&#160;</div><div class="line"><a name="l00197"></a><span class="lineno">  197</span>&#160;    val0 = (half8)((half4)(a0.s4, a1.s4, a2.s4, a3.s4), (half4)(a0.s5, a1.s5, a2.s5, a3.s5));</div><div class="line"><a name="l00198"></a><span class="lineno">  198</span>&#160;    vstore8(val0, 0, ((__global half *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 16);</div><div class="line"><a name="l00199"></a><span class="lineno">  199</span>&#160;</div><div class="line"><a name="l00200"></a><span class="lineno">  200</span>&#160;    val0 = (half8)((half4)(a0.s6, a1.s6, a2.s6, a3.s6), (half4)(a0.s7, a1.s7, a2.s7, a3.s7));</div><div class="line"><a name="l00201"></a><span class="lineno">  201</span>&#160;    vstore8(val0, 0, ((__global half *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 24);</div><div class="line"><a name="l00202"></a><span class="lineno">  202</span>&#160;}</div><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#l00201">helpers.h:201</a></div></div>
 462 <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>
 463 <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>
 464 <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>
 465 </div><!-- fragment -->
 466 </div>
 467 </div>
 468 <a class="anchor" id="acd423b6f992354e7c00137b20d687281"></a>
 469 <div class="memitem">
 470 <div class="memproto">
 471       <table class="memname">
 472         <tr>
 473           <td class="memname">__kernel void gemm_interleave4x4_f32 </td>
 474           <td>(</td>
 475           <td class="paramtype">__global uchar *&#160;</td>
 476           <td class="paramname"><em>src_ptr</em>, </td>
 477         </tr>
 478         <tr>
 479           <td class="paramkey"></td>
 480           <td></td>
 481           <td class="paramtype">uint&#160;</td>
 482           <td class="paramname"><em>src_stride_x</em>, </td>
 483         </tr>
 484         <tr>
 485           <td class="paramkey"></td>
 486           <td></td>
 487           <td class="paramtype">uint&#160;</td>
 488           <td class="paramname"><em>src_step_x</em>, </td>
 489         </tr>
 490         <tr>
 491           <td class="paramkey"></td>
 492           <td></td>
 493           <td class="paramtype">uint&#160;</td>
 494           <td class="paramname"><em>src_stride_y</em>, </td>
 495         </tr>
 496         <tr>
 497           <td class="paramkey"></td>
 498           <td></td>
 499           <td class="paramtype">uint&#160;</td>
 500           <td class="paramname"><em>src_step_y</em>, </td>
 501         </tr>
 502         <tr>
 503           <td class="paramkey"></td>
 504           <td></td>
 505           <td class="paramtype">uint&#160;</td>
 506           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 507         </tr>
 508         <tr>
 509           <td class="paramkey"></td>
 510           <td></td>
 511           <td class="paramtype">__global uchar *&#160;</td>
 512           <td class="paramname"><em>dst_ptr</em>, </td>
 513         </tr>
 514         <tr>
 515           <td class="paramkey"></td>
 516           <td></td>
 517           <td class="paramtype">uint&#160;</td>
 518           <td class="paramname"><em>dst_stride_x</em>, </td>
 519         </tr>
 520         <tr>
 521           <td class="paramkey"></td>
 522           <td></td>
 523           <td class="paramtype">uint&#160;</td>
 524           <td class="paramname"><em>dst_step_x</em>, </td>
 525         </tr>
 526         <tr>
 527           <td class="paramkey"></td>
 528           <td></td>
 529           <td class="paramtype">uint&#160;</td>
 530           <td class="paramname"><em>dst_stride_y</em>, </td>
 531         </tr>
 532         <tr>
 533           <td class="paramkey"></td>
 534           <td></td>
 535           <td class="paramtype">uint&#160;</td>
 536           <td class="paramname"><em>dst_step_y</em>, </td>
 537         </tr>
 538         <tr>
 539           <td class="paramkey"></td>
 540           <td></td>
 541           <td class="paramtype">uint&#160;</td>
 542           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
 543         </tr>
 544         <tr>
 545           <td></td>
 546           <td>)</td>
 547           <td></td><td></td>
 548         </tr>
 549       </table>
 550 </div><div class="memdoc">
 551
 552 <p>This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. </p>
 553 <dl class="params"><dt>Parameters</dt><dd>
 554   <table class="params">
 555     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: F32 </td></tr>
 556     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
 557     <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>
 558     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
 559     <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>
 560     <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 matrix </td></tr>
 561     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: F32 </td></tr>
 562     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
 563     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 564     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
 565     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 566     <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 matrix </td></tr>
 567   </table>
 568   </dd>
 569 </dl>
 570
 571 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00137">137</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
 572
 573 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00201">offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
 574 <div class="fragment"><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;    <span class="comment">/* Compute source and destination addresses */</span></div><div class="line"><a name="l00141"></a><span class="lineno">  141</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="l00142"></a><span class="lineno">  142</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> dst = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(dst);</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;    <span class="comment">/* Load values from Matrix A */</span></div><div class="line"><a name="l00145"></a><span class="lineno">  145</span>&#160;    float4 a0 = vload4(0, (__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0)));</div><div class="line"><a name="l00146"></a><span class="lineno">  146</span>&#160;    float4 a1 = vload4(0, (__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 1)));</div><div class="line"><a name="l00147"></a><span class="lineno">  147</span>&#160;    float4 a2 = vload4(0, (__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 2)));</div><div class="line"><a name="l00148"></a><span class="lineno">  148</span>&#160;    float4 a3 = vload4(0, (__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 3)));</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;    float4 val0 = (float4)(a0.s0, a1.s0, a2.s0, a3.s0);</div><div class="line"><a name="l00151"></a><span class="lineno">  151</span>&#160;    vstore4(val0, 0, ((__global <span class="keywordtype">float</span> *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 0);</div><div class="line"><a name="l00152"></a><span class="lineno">  152</span>&#160;</div><div class="line"><a name="l00153"></a><span class="lineno">  153</span>&#160;    val0 = (float4)(a0.s1, a1.s1, a2.s1, a3.s1);</div><div class="line"><a name="l00154"></a><span class="lineno">  154</span>&#160;    vstore4(val0, 0, ((__global <span class="keywordtype">float</span> *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 4);</div><div class="line"><a name="l00155"></a><span class="lineno">  155</span>&#160;</div><div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;    val0 = (float4)(a0.s2, a1.s2, a2.s2, a3.s2);</div><div class="line"><a name="l00157"></a><span class="lineno">  157</span>&#160;    vstore4(val0, 0, ((__global <span class="keywordtype">float</span> *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 8);</div><div class="line"><a name="l00158"></a><span class="lineno">  158</span>&#160;</div><div class="line"><a name="l00159"></a><span class="lineno">  159</span>&#160;    val0 = (float4)(a0.s3, a1.s3, a2.s3, a3.s3);</div><div class="line"><a name="l00160"></a><span class="lineno">  160</span>&#160;    vstore4(val0, 0, ((__global <span class="keywordtype">float</span> *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 12);</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;}</div><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#l00201">helpers.h:201</a></div></div>
 575 <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>
 576 <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>
 577 <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>
 578 </div><!-- fragment -->
 579 </div>
 580 </div>
 581 <a class="anchor" id="a830ba1cc0ad3c8cffa4a14424b2d0411"></a>
 582 <div class="memitem">
 583 <div class="memproto">
 584       <table class="memname">
 585         <tr>
 586           <td class="memname">__kernel void gemm_interleave4x4_u8 </td>
 587           <td>(</td>
 588           <td class="paramtype">__global uchar *&#160;</td>
 589           <td class="paramname"><em>src_ptr</em>, </td>
 590         </tr>
 591         <tr>
 592           <td class="paramkey"></td>
 593           <td></td>
 594           <td class="paramtype">uint&#160;</td>
 595           <td class="paramname"><em>src_stride_x</em>, </td>
 596         </tr>
 597         <tr>
 598           <td class="paramkey"></td>
 599           <td></td>
 600           <td class="paramtype">uint&#160;</td>
 601           <td class="paramname"><em>src_step_x</em>, </td>
 602         </tr>
 603         <tr>
 604           <td class="paramkey"></td>
 605           <td></td>
 606           <td class="paramtype">uint&#160;</td>
 607           <td class="paramname"><em>src_stride_y</em>, </td>
 608         </tr>
 609         <tr>
 610           <td class="paramkey"></td>
 611           <td></td>
 612           <td class="paramtype">uint&#160;</td>
 613           <td class="paramname"><em>src_step_y</em>, </td>
 614         </tr>
 615         <tr>
 616           <td class="paramkey"></td>
 617           <td></td>
 618           <td class="paramtype">uint&#160;</td>
 619           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 620         </tr>
 621         <tr>
 622           <td class="paramkey"></td>
 623           <td></td>
 624           <td class="paramtype">__global uchar *&#160;</td>
 625           <td class="paramname"><em>dst_ptr</em>, </td>
 626         </tr>
 627         <tr>
 628           <td class="paramkey"></td>
 629           <td></td>
 630           <td class="paramtype">uint&#160;</td>
 631           <td class="paramname"><em>dst_stride_x</em>, </td>
 632         </tr>
 633         <tr>
 634           <td class="paramkey"></td>
 635           <td></td>
 636           <td class="paramtype">uint&#160;</td>
 637           <td class="paramname"><em>dst_step_x</em>, </td>
 638         </tr>
 639         <tr>
 640           <td class="paramkey"></td>
 641           <td></td>
 642           <td class="paramtype">uint&#160;</td>
 643           <td class="paramname"><em>dst_stride_y</em>, </td>
 644         </tr>
 645         <tr>
 646           <td class="paramkey"></td>
 647           <td></td>
 648           <td class="paramtype">uint&#160;</td>
 649           <td class="paramname"><em>dst_step_y</em>, </td>
 650         </tr>
 651         <tr>
 652           <td class="paramkey"></td>
 653           <td></td>
 654           <td class="paramtype">uint&#160;</td>
 655           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
 656         </tr>
 657         <tr>
 658           <td></td>
 659           <td>)</td>
 660           <td></td><td></td>
 661         </tr>
 662       </table>
 663 </div><div class="memdoc">
 664
 665 <p>This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. </p>
 666 <dl class="params"><dt>Parameters</dt><dd>
 667   <table class="params">
 668     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U8 </td></tr>
 669     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
 670     <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>
 671     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
 672     <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>
 673     <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 matrix </td></tr>
 674     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: U8 </td></tr>
 675     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
 676     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 677     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
 678     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 679     <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 matrix </td></tr>
 680   </table>
 681   </dd>
 682 </dl>
 683
 684 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00219">219</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
 685
 686 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00201">offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
 687 <div class="fragment"><div class="line"><a name="l00221"></a><span class="lineno">  221</span>&#160;{</div><div class="line"><a name="l00222"></a><span class="lineno">  222</span>&#160;    <span class="comment">/* Compute source and destination addresses */</span></div><div class="line"><a name="l00223"></a><span class="lineno">  223</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="l00224"></a><span class="lineno">  224</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> dst = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00225"></a><span class="lineno">  225</span>&#160;</div><div class="line"><a name="l00226"></a><span class="lineno">  226</span>&#160;    <span class="comment">/* Load values from Matrix A */</span></div><div class="line"><a name="l00227"></a><span class="lineno">  227</span>&#160;    uchar16 a0 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0)));</div><div class="line"><a name="l00228"></a><span class="lineno">  228</span>&#160;    uchar16 a1 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 1)));</div><div class="line"><a name="l00229"></a><span class="lineno">  229</span>&#160;    uchar16 a2 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 2)));</div><div class="line"><a name="l00230"></a><span class="lineno">  230</span>&#160;    uchar16 a3 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 3)));</div><div class="line"><a name="l00231"></a><span class="lineno">  231</span>&#160;</div><div class="line"><a name="l00232"></a><span class="lineno">  232</span>&#160;    uchar16 val0 = (uchar16)((uchar4)(a0.s0, a1.s0, a2.s0, a3.s0), (uchar4)(a0.s1, a1.s1, a2.s1, a3.s1),</div><div class="line"><a name="l00233"></a><span class="lineno">  233</span>&#160;                             (uchar4)(a0.s2, a1.s2, a2.s2, a3.s2), (uchar4)(a0.s3, a1.s3, a2.s3, a3.s3));</div><div class="line"><a name="l00234"></a><span class="lineno">  234</span>&#160;    vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 0);</div><div class="line"><a name="l00235"></a><span class="lineno">  235</span>&#160;</div><div class="line"><a name="l00236"></a><span class="lineno">  236</span>&#160;    val0 = (uchar16)((uchar4)(a0.s4, a1.s4, a2.s4, a3.s4), (uchar4)(a0.s5, a1.s5, a2.s5, a3.s5),</div><div class="line"><a name="l00237"></a><span class="lineno">  237</span>&#160;                     (uchar4)(a0.s6, a1.s6, a2.s6, a3.s6), (uchar4)(a0.s7, a1.s7, a2.s7, a3.s7));</div><div class="line"><a name="l00238"></a><span class="lineno">  238</span>&#160;    vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 16);</div><div class="line"><a name="l00239"></a><span class="lineno">  239</span>&#160;</div><div class="line"><a name="l00240"></a><span class="lineno">  240</span>&#160;    val0 = (uchar16)((uchar4)(a0.s8, a1.s8, a2.s8, a3.s8), (uchar4)(a0.s9, a1.s9, a2.s9, a3.s9),</div><div class="line"><a name="l00241"></a><span class="lineno">  241</span>&#160;                     (uchar4)(a0.sA, a1.sA, a2.sA, a3.sA), (uchar4)(a0.sB, a1.sB, a2.sB, a3.sB));</div><div class="line"><a name="l00242"></a><span class="lineno">  242</span>&#160;    vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 32);</div><div class="line"><a name="l00243"></a><span class="lineno">  243</span>&#160;</div><div class="line"><a name="l00244"></a><span class="lineno">  244</span>&#160;    val0 = (uchar16)((uchar4)(a0.sC, a1.sC, a2.sC, a3.sC), (uchar4)(a0.sD, a1.sD, a2.sD, a3.sD),</div><div class="line"><a name="l00245"></a><span class="lineno">  245</span>&#160;                     (uchar4)(a0.sE, a1.sE, a2.sE, a3.sE), (uchar4)(a0.sF, a1.sF, a2.sF, a3.sF));</div><div class="line"><a name="l00246"></a><span class="lineno">  246</span>&#160;    vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 48);</div><div class="line"><a name="l00247"></a><span class="lineno">  247</span>&#160;}</div><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#l00201">helpers.h:201</a></div></div>
 688 <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>
 689 <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>
 690 <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>
 691 </div><!-- fragment -->
 692 </div>
 693 </div>
 694 <a class="anchor" id="a675879eb72964ddbd5f71155da3b0cc3"></a>
 695 <div class="memitem">
 696 <div class="memproto">
 697       <table class="memname">
 698         <tr>
 699           <td class="memname">__kernel void gemm_transpose1x16_u8 </td>
 700           <td>(</td>
 701           <td class="paramtype">__global uchar *&#160;</td>
 702           <td class="paramname"><em>src_ptr</em>, </td>
 703         </tr>
 704         <tr>
 705           <td class="paramkey"></td>
 706           <td></td>
 707           <td class="paramtype">uint&#160;</td>
 708           <td class="paramname"><em>src_stride_x</em>, </td>
 709         </tr>
 710         <tr>
 711           <td class="paramkey"></td>
 712           <td></td>
 713           <td class="paramtype">uint&#160;</td>
 714           <td class="paramname"><em>src_step_x</em>, </td>
 715         </tr>
 716         <tr>
 717           <td class="paramkey"></td>
 718           <td></td>
 719           <td class="paramtype">uint&#160;</td>
 720           <td class="paramname"><em>src_stride_y</em>, </td>
 721         </tr>
 722         <tr>
 723           <td class="paramkey"></td>
 724           <td></td>
 725           <td class="paramtype">uint&#160;</td>
 726           <td class="paramname"><em>src_step_y</em>, </td>
 727         </tr>
 728         <tr>
 729           <td class="paramkey"></td>
 730           <td></td>
 731           <td class="paramtype">uint&#160;</td>
 732           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 733         </tr>
 734         <tr>
 735           <td class="paramkey"></td>
 736           <td></td>
 737           <td class="paramtype">__global uchar *&#160;</td>
 738           <td class="paramname"><em>dst_ptr</em>, </td>
 739         </tr>
 740         <tr>
 741           <td class="paramkey"></td>
 742           <td></td>
 743           <td class="paramtype">uint&#160;</td>
 744           <td class="paramname"><em>dst_stride_x</em>, </td>
 745         </tr>
 746         <tr>
 747           <td class="paramkey"></td>
 748           <td></td>
 749           <td class="paramtype">uint&#160;</td>
 750           <td class="paramname"><em>dst_step_x</em>, </td>
 751         </tr>
 752         <tr>
 753           <td class="paramkey"></td>
 754           <td></td>
 755           <td class="paramtype">uint&#160;</td>
 756           <td class="paramname"><em>dst_stride_y</em>, </td>
 757         </tr>
 758         <tr>
 759           <td class="paramkey"></td>
 760           <td></td>
 761           <td class="paramtype">uint&#160;</td>
 762           <td class="paramname"><em>dst_step_y</em>, </td>
 763         </tr>
 764         <tr>
 765           <td class="paramkey"></td>
 766           <td></td>
 767           <td class="paramtype">uint&#160;</td>
 768           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
 769         </tr>
 770         <tr>
 771           <td></td>
 772           <td>)</td>
 773           <td></td><td></td>
 774         </tr>
 775       </table>
 776 </div><div class="memdoc">
 777
 778 <p>This OpenCL kernel computes the "vector" 1x16 transposition of input matrix. </p>
 779 <dl class="params"><dt>Parameters</dt><dd>
 780   <table class="params">
 781     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U8 </td></tr>
 782     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
 783     <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>
 784     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
 785     <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>
 786     <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 matrix </td></tr>
 787     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: U8 </td></tr>
 788     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
 789     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 790     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
 791     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 792     <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 matrix </td></tr>
 793   </table>
 794   </dd>
 795 </dl>
 796
 797 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00105">105</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
 798
 799 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
 800 <div class="fragment"><div class="line"><a name="l00107"></a><span class="lineno">  107</span>&#160;{</div><div class="line"><a name="l00108"></a><span class="lineno">  108</span>&#160;    uint x = get_global_id(0);</div><div class="line"><a name="l00109"></a><span class="lineno">  109</span>&#160;    uint y = get_global_id(1);</div><div class="line"><a name="l00110"></a><span class="lineno">  110</span>&#160;</div><div class="line"><a name="l00111"></a><span class="lineno">  111</span>&#160;    <span class="comment">/* Compute address for Matrix B - source */</span></div><div class="line"><a name="l00112"></a><span class="lineno">  112</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="l00113"></a><span class="lineno">  113</span>&#160;</div><div class="line"><a name="l00114"></a><span class="lineno">  114</span>&#160;    <span class="comment">/* Compute address for Matrix B transposed - destination. X and Y are swapped */</span></div><div class="line"><a name="l00115"></a><span class="lineno">  115</span>&#160;    uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));</div><div class="line"><a name="l00116"></a><span class="lineno">  116</span>&#160;</div><div class="line"><a name="l00117"></a><span class="lineno">  117</span>&#160;    uchar16 b0 = vload16(0, (__global uchar *)src.ptr);</div><div class="line"><a name="l00118"></a><span class="lineno">  118</span>&#160;</div><div class="line"><a name="l00119"></a><span class="lineno">  119</span>&#160;    vstore16(b0, 0, (__global uchar *)(dst_ptr + dst_addr_in_bytes));</div><div class="line"><a name="l00120"></a><span class="lineno">  120</span>&#160;}</div><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>
 801 <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>
 802 </div><!-- fragment -->
 803 </div>
 804 </div>
 805 <a class="anchor" id="a21d7061f36aec78824e368f2ef1cafc1"></a>
 806 <div class="memitem">
 807 <div class="memproto">
 808       <table class="memname">
 809         <tr>
 810           <td class="memname">__kernel void gemm_transpose1x4_f32 </td>
 811           <td>(</td>
 812           <td class="paramtype">__global uchar *&#160;</td>
 813           <td class="paramname"><em>src_ptr</em>, </td>
 814         </tr>
 815         <tr>
 816           <td class="paramkey"></td>
 817           <td></td>
 818           <td class="paramtype">uint&#160;</td>
 819           <td class="paramname"><em>src_stride_x</em>, </td>
 820         </tr>
 821         <tr>
 822           <td class="paramkey"></td>
 823           <td></td>
 824           <td class="paramtype">uint&#160;</td>
 825           <td class="paramname"><em>src_step_x</em>, </td>
 826         </tr>
 827         <tr>
 828           <td class="paramkey"></td>
 829           <td></td>
 830           <td class="paramtype">uint&#160;</td>
 831           <td class="paramname"><em>src_stride_y</em>, </td>
 832         </tr>
 833         <tr>
 834           <td class="paramkey"></td>
 835           <td></td>
 836           <td class="paramtype">uint&#160;</td>
 837           <td class="paramname"><em>src_step_y</em>, </td>
 838         </tr>
 839         <tr>
 840           <td class="paramkey"></td>
 841           <td></td>
 842           <td class="paramtype">uint&#160;</td>
 843           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 844         </tr>
 845         <tr>
 846           <td class="paramkey"></td>
 847           <td></td>
 848           <td class="paramtype">__global uchar *&#160;</td>
 849           <td class="paramname"><em>dst_ptr</em>, </td>
 850         </tr>
 851         <tr>
 852           <td class="paramkey"></td>
 853           <td></td>
 854           <td class="paramtype">uint&#160;</td>
 855           <td class="paramname"><em>dst_stride_x</em>, </td>
 856         </tr>
 857         <tr>
 858           <td class="paramkey"></td>
 859           <td></td>
 860           <td class="paramtype">uint&#160;</td>
 861           <td class="paramname"><em>dst_step_x</em>, </td>
 862         </tr>
 863         <tr>
 864           <td class="paramkey"></td>
 865           <td></td>
 866           <td class="paramtype">uint&#160;</td>
 867           <td class="paramname"><em>dst_stride_y</em>, </td>
 868         </tr>
 869         <tr>
 870           <td class="paramkey"></td>
 871           <td></td>
 872           <td class="paramtype">uint&#160;</td>
 873           <td class="paramname"><em>dst_step_y</em>, </td>
 874         </tr>
 875         <tr>
 876           <td class="paramkey"></td>
 877           <td></td>
 878           <td class="paramtype">uint&#160;</td>
 879           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
 880         </tr>
 881         <tr>
 882           <td></td>
 883           <td>)</td>
 884           <td></td><td></td>
 885         </tr>
 886       </table>
 887 </div><div class="memdoc">
 888
 889 <p>This OpenCL kernel computes the "vector" 1x4 transposition of input matrix. </p>
 890 <dl class="params"><dt>Parameters</dt><dd>
 891   <table class="params">
 892     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: F32 </td></tr>
 893     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
 894     <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>
 895     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
 896     <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>
 897     <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 matrix </td></tr>
 898     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: F32 </td></tr>
 899     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
 900     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 901     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
 902     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 903     <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 matrix </td></tr>
 904   </table>
 905   </dd>
 906 </dl>
 907
 908 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00041">41</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
 909
 910 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
 911 <div class="fragment"><div class="line"><a name="l00043"></a><span class="lineno">   43</span>&#160;{</div><div class="line"><a name="l00044"></a><span class="lineno">   44</span>&#160;    uint x = get_global_id(0);</div><div class="line"><a name="l00045"></a><span class="lineno">   45</span>&#160;    uint y = get_global_id(1);</div><div class="line"><a name="l00046"></a><span class="lineno">   46</span>&#160;</div><div class="line"><a name="l00047"></a><span class="lineno">   47</span>&#160;    <span class="comment">/* Compute address for Matrix B - source */</span></div><div class="line"><a name="l00048"></a><span class="lineno">   48</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="l00049"></a><span class="lineno">   49</span>&#160;</div><div class="line"><a name="l00050"></a><span class="lineno">   50</span>&#160;    <span class="comment">/* Compute address for Matrix B transposed - destination. X and Y are swapped */</span></div><div class="line"><a name="l00051"></a><span class="lineno">   51</span>&#160;    uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));</div><div class="line"><a name="l00052"></a><span class="lineno">   52</span>&#160;</div><div class="line"><a name="l00053"></a><span class="lineno">   53</span>&#160;    float4 b0 = vload4(0, (__global <span class="keywordtype">float</span> *)src.ptr);</div><div class="line"><a name="l00054"></a><span class="lineno">   54</span>&#160;</div><div class="line"><a name="l00055"></a><span class="lineno">   55</span>&#160;    vstore4(b0, 0, (__global <span class="keywordtype">float</span> *)(dst_ptr + dst_addr_in_bytes));</div><div class="line"><a name="l00056"></a><span class="lineno">   56</span>&#160;}</div><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>
 912 <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>
 913 </div><!-- fragment -->
 914 </div>
 915 </div>
 916 <a class="anchor" id="a11d87ecaf1388b4f57404e062265b374"></a>
 917 <div class="memitem">
 918 <div class="memproto">
 919       <table class="memname">
 920         <tr>
 921           <td class="memname">__kernel void gemm_transpose1x8_f16 </td>
 922           <td>(</td>
 923           <td class="paramtype">__global uchar *&#160;</td>
 924           <td class="paramname"><em>src_ptr</em>, </td>
 925         </tr>
 926         <tr>
 927           <td class="paramkey"></td>
 928           <td></td>
 929           <td class="paramtype">uint&#160;</td>
 930           <td class="paramname"><em>src_stride_x</em>, </td>
 931         </tr>
 932         <tr>
 933           <td class="paramkey"></td>
 934           <td></td>
 935           <td class="paramtype">uint&#160;</td>
 936           <td class="paramname"><em>src_step_x</em>, </td>
 937         </tr>
 938         <tr>
 939           <td class="paramkey"></td>
 940           <td></td>
 941           <td class="paramtype">uint&#160;</td>
 942           <td class="paramname"><em>src_stride_y</em>, </td>
 943         </tr>
 944         <tr>
 945           <td class="paramkey"></td>
 946           <td></td>
 947           <td class="paramtype">uint&#160;</td>
 948           <td class="paramname"><em>src_step_y</em>, </td>
 949         </tr>
 950         <tr>
 951           <td class="paramkey"></td>
 952           <td></td>
 953           <td class="paramtype">uint&#160;</td>
 954           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 955         </tr>
 956         <tr>
 957           <td class="paramkey"></td>
 958           <td></td>
 959           <td class="paramtype">__global uchar *&#160;</td>
 960           <td class="paramname"><em>dst_ptr</em>, </td>
 961         </tr>
 962         <tr>
 963           <td class="paramkey"></td>
 964           <td></td>
 965           <td class="paramtype">uint&#160;</td>
 966           <td class="paramname"><em>dst_stride_x</em>, </td>
 967         </tr>
 968         <tr>
 969           <td class="paramkey"></td>
 970           <td></td>
 971           <td class="paramtype">uint&#160;</td>
 972           <td class="paramname"><em>dst_step_x</em>, </td>
 973         </tr>
 974         <tr>
 975           <td class="paramkey"></td>
 976           <td></td>
 977           <td class="paramtype">uint&#160;</td>
 978           <td class="paramname"><em>dst_stride_y</em>, </td>
 979         </tr>
 980         <tr>
 981           <td class="paramkey"></td>
 982           <td></td>
 983           <td class="paramtype">uint&#160;</td>
 984           <td class="paramname"><em>dst_step_y</em>, </td>
 985         </tr>
 986         <tr>
 987           <td class="paramkey"></td>
 988           <td></td>
 989           <td class="paramtype">uint&#160;</td>
 990           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
 991         </tr>
 992         <tr>
 993           <td></td>
 994           <td>)</td>
 995           <td></td><td></td>
 996         </tr>
 997       </table>
 998 </div><div class="memdoc">
 999
1000 <p>This OpenCL kernel computes the "vector" 1x8 transposition of input matrix. </p>
1001 <dl class="params"><dt>Parameters</dt><dd>
1002   <table class="params">
1003     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: F16 </td></tr>
1004     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
1005     <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>
1006     <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
1007     <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>
1008     <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 matrix </td></tr>
1009     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: F16 </td></tr>
1010     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
1011     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
1012     <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
1013     <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
1014     <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 matrix </td></tr>
1015   </table>
1016   </dd>
1017 </dl>
1018
1019 <p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00073">73</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
1020
1021 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, and <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>.</p>
1022 <div class="fragment"><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;    uint x = get_global_id(0);</div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;    uint y = get_global_id(1);</div><div class="line"><a name="l00078"></a><span class="lineno">   78</span>&#160;</div><div class="line"><a name="l00079"></a><span class="lineno">   79</span>&#160;    <span class="comment">/* Compute address for Matrix B - source */</span></div><div class="line"><a name="l00080"></a><span class="lineno">   80</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="l00081"></a><span class="lineno">   81</span>&#160;</div><div class="line"><a name="l00082"></a><span class="lineno">   82</span>&#160;    <span class="comment">/* Compute address for Matrix B transposed - destination. X and Y are swapped */</span></div><div class="line"><a name="l00083"></a><span class="lineno">   83</span>&#160;    uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));</div><div class="line"><a name="l00084"></a><span class="lineno">   84</span>&#160;</div><div class="line"><a name="l00085"></a><span class="lineno">   85</span>&#160;    half8 b0 = vload8(0, (__global half *)src.ptr);</div><div class="line"><a name="l00086"></a><span class="lineno">   86</span>&#160;</div><div class="line"><a name="l00087"></a><span class="lineno">   87</span>&#160;    vstore8(b0, 0, (__global half *)(dst_ptr + dst_addr_in_bytes));</div><div class="line"><a name="l00088"></a><span class="lineno">   88</span>&#160;}</div><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>
1023 <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>
1024 </div><!-- fragment -->
1025 </div>
1026 </div>
1027 </div><!-- contents -->
1028 </div><!-- doc-content -->
1029 <!-- start footer part -->
1030 <div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
1031   <ul>
1032     <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="gemm_8cl.xhtml">gemm.cl</a></li>
1033     <li class="footer">Generated on Wed Apr 12 2017 14:26:05 for ARM Compute Library by
1034     <a href="http://www.doxygen.org/index.html">
1035     <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
1036   </ul>
1037 </div>
1038 </body>
1039 </html>