documentation/softmax__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/softmax_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('softmax__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="#define-members">Macros</a> &#124;
 117 <a href="#func-members">Functions</a> &#124;
 118 <a href="#var-members">Variables</a>  </div>
 119   <div class="headertitle">
 120 <div class="title">softmax_layer.cl File Reference</div>  </div>
 121 </div><!--header-->
 122 <div class="contents">
 123 <div class="textblock"><code>#include &quot;<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
 124 </div>
 125 <p><a href="softmax__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
 126 <table class="memberdecls">
 127 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
 128 Macros</h2></td></tr>
 129 <tr class="memitem:a80b22c555ddadb47cc6ca338a9c49126"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a>&#160;&#160;&#160;FLT_MIN</td></tr>
 130 <tr class="separator:a80b22c555ddadb47cc6ca338a9c49126"><td class="memSeparator" colspan="2">&#160;</td></tr>
 131 <tr class="memitem:af5b2e33e3c5fcaab3a213f26c2300170"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>&#160;&#160;&#160;int</td></tr>
 132 <tr class="separator:af5b2e33e3c5fcaab3a213f26c2300170"><td class="memSeparator" colspan="2">&#160;</td></tr>
 133 <tr class="memitem:afb8c72ce35c4a1f4a2588d6573e54aa1"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>&#160;&#160;&#160;float</td></tr>
 134 <tr class="separator:afb8c72ce35c4a1f4a2588d6573e54aa1"><td class="memSeparator" colspan="2">&#160;</td></tr>
 135 </table><table class="memberdecls">
 136 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
 137 Functions</h2></td></tr>
 138 <tr class="memitem:a67da577562e431fc20c5cd8a1e8ddf21"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a67da577562e431fc20c5cd8a1e8ddf21">softmax_layer_max</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, uint width)</td></tr>
 139 <tr class="memdesc:a67da577562e431fc20c5cd8a1e8ddf21"><td class="mdescLeft">&#160;</td><td class="mdescRight">Identifies the maximum value across the 1st dimension.  <a href="#a67da577562e431fc20c5cd8a1e8ddf21">More...</a><br /></td></tr>
 140 <tr class="separator:a67da577562e431fc20c5cd8a1e8ddf21"><td class="memSeparator" colspan="2">&#160;</td></tr>
 141 <tr class="memitem:aaf2858ffa79555d18154eae8a32db43e"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#aaf2858ffa79555d18154eae8a32db43e">softmax_layer_shift_exp_sum</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 *max_ptr, uint max_stride_x, uint max_step_x, uint max_stride_y, uint max_step_y, uint max_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 *sum_ptr, uint sum_stride_x, uint sum_step_x, uint sum_stride_y, uint sum_step_y, uint sum_offset_first_element_in_bytes, uint width)</td></tr>
 142 <tr class="memdesc:aaf2858ffa79555d18154eae8a32db43e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Shifts the values of the input tensor by the max calculated in softmax_layer_max kernel, then gets the exponent of each element as sums all elements across each row.  <a href="#aaf2858ffa79555d18154eae8a32db43e">More...</a><br /></td></tr>
 143 <tr class="separator:aaf2858ffa79555d18154eae8a32db43e"><td class="memSeparator" colspan="2">&#160;</td></tr>
 144 <tr class="memitem:a317069d3b4107b4b8157c8b09e30745f"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a317069d3b4107b4b8157c8b09e30745f">softmax_layer_norm</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 *sum_ptr, uint sum_stride_x, uint sum_step_x, uint sum_stride_y, uint sum_step_y, uint sum_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>
 145 <tr class="memdesc:a317069d3b4107b4b8157c8b09e30745f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel.  <a href="#a317069d3b4107b4b8157c8b09e30745f">More...</a><br /></td></tr>
 146 <tr class="separator:a317069d3b4107b4b8157c8b09e30745f"><td class="memSeparator" colspan="2">&#160;</td></tr>
 147 </table><table class="memberdecls">
 148 <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="var-members"></a>
 149 Variables</h2></td></tr>
 150 <tr class="memitem:a5505428916dfe70b40b32686c57e8436"><td class="memItemLeft" align="right" valign="top">__constant float16&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a5505428916dfe70b40b32686c57e8436">type_min</a> = ( float16 )( FLT_MIN )</td></tr>
 151 <tr class="separator:a5505428916dfe70b40b32686c57e8436"><td class="memSeparator" colspan="2">&#160;</td></tr>
 152 <tr class="memitem:a0712735973f172ac9efc7d48a31e47ad"><td class="memItemLeft" align="right" valign="top">__constant uint16&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a> = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td></tr>
 153 <tr class="separator:a0712735973f172ac9efc7d48a31e47ad"><td class="memSeparator" colspan="2">&#160;</td></tr>
 154 </table>
 155 <h2 class="groupheader">Macro Definition Documentation</h2>
 156 <a class="anchor" id="afb8c72ce35c4a1f4a2588d6573e54aa1"></a>
 157 <div class="memitem">
 158 <div class="memproto">
 159       <table class="memname">
 160         <tr>
 161           <td class="memname">#define DATA_TYPE&#160;&#160;&#160;float</td>
 162         </tr>
 163       </table>
 164 </div><div class="memdoc">
 165
 166 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00033">33</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 167
 168 <p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00207">softmax_layer_norm()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00132">softmax_layer_shift_exp_sum()</a>.</p>
 169
 170 </div>
 171 </div>
 172 <a class="anchor" id="a80b22c555ddadb47cc6ca338a9c49126"></a>
 173 <div class="memitem">
 174 <div class="memproto">
 175       <table class="memname">
 176         <tr>
 177           <td class="memname">#define MINVAL&#160;&#160;&#160;FLT_MIN</td>
 178         </tr>
 179       </table>
 180 </div><div class="memdoc">
 181
 182 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00031">31</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 183
 184 </div>
 185 </div>
 186 <a class="anchor" id="af5b2e33e3c5fcaab3a213f26c2300170"></a>
 187 <div class="memitem">
 188 <div class="memproto">
 189       <table class="memname">
 190         <tr>
 191           <td class="memname">#define SELECT_DATA_TYPE&#160;&#160;&#160;int</td>
 192         </tr>
 193       </table>
 194 </div><div class="memdoc">
 195
 196 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00032">32</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 197
 198 <p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00059">softmax_layer_max()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00132">softmax_layer_shift_exp_sum()</a>.</p>
 199
 200 </div>
 201 </div>
 202 <h2 class="groupheader">Function Documentation</h2>
 203 <a class="anchor" id="a67da577562e431fc20c5cd8a1e8ddf21"></a>
 204 <div class="memitem">
 205 <div class="memproto">
 206       <table class="memname">
 207         <tr>
 208           <td class="memname">__kernel void softmax_layer_max </td>
 209           <td>(</td>
 210           <td class="paramtype">__global uchar *&#160;</td>
 211           <td class="paramname"><em>src_ptr</em>, </td>
 212         </tr>
 213         <tr>
 214           <td class="paramkey"></td>
 215           <td></td>
 216           <td class="paramtype">uint&#160;</td>
 217           <td class="paramname"><em>src_stride_x</em>, </td>
 218         </tr>
 219         <tr>
 220           <td class="paramkey"></td>
 221           <td></td>
 222           <td class="paramtype">uint&#160;</td>
 223           <td class="paramname"><em>src_step_x</em>, </td>
 224         </tr>
 225         <tr>
 226           <td class="paramkey"></td>
 227           <td></td>
 228           <td class="paramtype">uint&#160;</td>
 229           <td class="paramname"><em>src_stride_y</em>, </td>
 230         </tr>
 231         <tr>
 232           <td class="paramkey"></td>
 233           <td></td>
 234           <td class="paramtype">uint&#160;</td>
 235           <td class="paramname"><em>src_step_y</em>, </td>
 236         </tr>
 237         <tr>
 238           <td class="paramkey"></td>
 239           <td></td>
 240           <td class="paramtype">uint&#160;</td>
 241           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 242         </tr>
 243         <tr>
 244           <td class="paramkey"></td>
 245           <td></td>
 246           <td class="paramtype">__global uchar *&#160;</td>
 247           <td class="paramname"><em>dst_ptr</em>, </td>
 248         </tr>
 249         <tr>
 250           <td class="paramkey"></td>
 251           <td></td>
 252           <td class="paramtype">uint&#160;</td>
 253           <td class="paramname"><em>dst_stride_x</em>, </td>
 254         </tr>
 255         <tr>
 256           <td class="paramkey"></td>
 257           <td></td>
 258           <td class="paramtype">uint&#160;</td>
 259           <td class="paramname"><em>dst_step_x</em>, </td>
 260         </tr>
 261         <tr>
 262           <td class="paramkey"></td>
 263           <td></td>
 264           <td class="paramtype">uint&#160;</td>
 265           <td class="paramname"><em>dst_stride_y</em>, </td>
 266         </tr>
 267         <tr>
 268           <td class="paramkey"></td>
 269           <td></td>
 270           <td class="paramtype">uint&#160;</td>
 271           <td class="paramname"><em>dst_step_y</em>, </td>
 272         </tr>
 273         <tr>
 274           <td class="paramkey"></td>
 275           <td></td>
 276           <td class="paramtype">uint&#160;</td>
 277           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
 278         </tr>
 279         <tr>
 280           <td class="paramkey"></td>
 281           <td></td>
 282           <td class="paramtype">uint&#160;</td>
 283           <td class="paramname"><em>width</em>&#160;</td>
 284         </tr>
 285         <tr>
 286           <td></td>
 287           <td>)</td>
 288           <td></td><td></td>
 289         </tr>
 290       </table>
 291 </div><div class="memdoc">
 292
 293 <p>Identifies the maximum value across the 1st dimension. </p>
 294 <dl class="section note"><dt>Note</dt><dd>Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short </dd>
 295 <dd>
 296 In case F16 is used -DUSE_HALF must be passed otherwise the kernel will default to used F32. </dd>
 297 <dd>
 298 In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed.</dd></dl>
 299 <dl class="params"><dt>Parameters</dt><dd>
 300   <table class="params">
 301     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16, F32 </td></tr>
 302     <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>
 303     <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>
 304     <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>
 305     <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>
 306     <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>
 307     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor slice. Supported data types: F16, F32 </td></tr>
 308     <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>
 309     <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>
 310     <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>
 311     <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>
 312     <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>
 313     <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
 314   </table>
 315   </dd>
 316 </dl>
 317
 318 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00059">59</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 319
 320 <p>References <a class="el" href="helpers_8h_source.xhtml#l00035">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00037">idx16</a>, <a class="el" href="helpers_8h_source.xhtml#l00201">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00032">SELECT_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00036">type_min</a>, and <a class="el" href="helpers_8h_source.xhtml#l00032">VEC_DATA_TYPE</a>.</p>
 321 <div class="fragment"><div class="line"><a name="l00063"></a><span class="lineno">   63</span>&#160;{</div><div class="line"><a name="l00064"></a><span class="lineno">   64</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="l00065"></a><span class="lineno">   65</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="l00066"></a><span class="lineno">   66</span>&#160;</div><div class="line"><a name="l00067"></a><span class="lineno">   67</span>&#160;    <span class="comment">// Initialize local maximum</span></div><div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;    max_val = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16))<a class="code" href="softmax__layer_8cl.xhtml#a5505428916dfe70b40b32686c57e8436">type_min</a>;</div><div class="line"><a name="l00070"></a><span class="lineno">   70</span>&#160;</div><div class="line"><a name="l00071"></a><span class="lineno">   71</span>&#160;    <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00072"></a><span class="lineno">   72</span>&#160;    const uint width4 = width &gt;&gt; 4;</div><div class="line"><a name="l00073"></a><span class="lineno">   73</span>&#160;    for(uint i = 0; i &lt; width4; i++)</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;        <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00076"></a><span class="lineno">   76</span>&#160;        data    = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; 4, 0));</div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;        max_val = max(data, max_val);</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;</div><div class="line"><a name="l00080"></a><span class="lineno">   80</span>&#160;<span class="preprocessor">#if defined NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00081"></a><span class="lineno">   81</span>&#160;    <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00082"></a><span class="lineno">   82</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00083"></a><span class="lineno">   83</span>&#160;    data = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width4 &lt;&lt; 4, 0));</div><div class="line"><a name="l00084"></a><span class="lineno">   84</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 16)</div><div class="line"><a name="l00085"></a><span class="lineno">   85</span>&#160;    widx    = <a class="code" href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint16)(width4 &lt;&lt; 4) + <a class="code" href="softmax__layer_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a>) &lt; width, <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 16));</div><div class="line"><a name="l00086"></a><span class="lineno">   86</span>&#160;    max_val = max(max_val, select(type_min, data, widx));</div><div class="line"><a name="l00087"></a><span class="lineno">   87</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00088"></a><span class="lineno">   88</span>&#160;</div><div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00090"></a><span class="lineno">   90</span>&#160;    max_val.s01234567 = max(max_val.s01234567, max_val.s89ABCDEF);</div><div class="line"><a name="l00091"></a><span class="lineno">   91</span>&#160;    max_val.s0123     = max(max_val.s0123, max_val.s4567);</div><div class="line"><a name="l00092"></a><span class="lineno">   92</span>&#160;    max_val.s01       = max(max_val.s01, max_val.s23);</div><div class="line"><a name="l00093"></a><span class="lineno">   93</span>&#160;    max_val.s0        = max(max_val.s0, max_val.s1);</div><div class="line"><a name="l00094"></a><span class="lineno">   94</span>&#160;</div><div class="line"><a name="l00095"></a><span class="lineno">   95</span>&#160;    <span class="comment">// Store result</span></div><div class="line"><a name="l00096"></a><span class="lineno">   96</span>&#160;    *((__global DATA_TYPE *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = max_val.s0;</div><div class="line"><a name="l00097"></a><span class="lineno">   97</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>
 322 <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>
 323 <div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00032">helpers.h:32</a></div></div>
 324 <div class="ttc" id="softmax__layer_8cl_xhtml_af5b2e33e3c5fcaab3a213f26c2300170"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a></div><div class="ttdeci">#define SELECT_DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00032">softmax_layer.cl:32</a></div></div>
 325 <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>
 326 <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>
 327 <div class="ttc" id="softmax__layer_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00033">softmax_layer.cl:33</a></div></div>
 328 <div class="ttc" id="softmax__layer_8cl_xhtml_a5505428916dfe70b40b32686c57e8436"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a5505428916dfe70b40b32686c57e8436">type_min</a></div><div class="ttdeci">__constant float16 type_min</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00036">softmax_layer.cl:36</a></div></div>
 329 <div class="ttc" id="softmax__layer_8cl_xhtml_a0712735973f172ac9efc7d48a31e47ad"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a></div><div class="ttdeci">__constant uint16 idx16</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00037">softmax_layer.cl:37</a></div></div>
 330 <div class="ttc" id="helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00035">helpers.h:35</a></div></div>
 331 </div><!-- fragment -->
 332 </div>
 333 </div>
 334 <a class="anchor" id="a317069d3b4107b4b8157c8b09e30745f"></a>
 335 <div class="memitem">
 336 <div class="memproto">
 337       <table class="memname">
 338         <tr>
 339           <td class="memname">__kernel void softmax_layer_norm </td>
 340           <td>(</td>
 341           <td class="paramtype">__global uchar *&#160;</td>
 342           <td class="paramname"><em>src_ptr</em>, </td>
 343         </tr>
 344         <tr>
 345           <td class="paramkey"></td>
 346           <td></td>
 347           <td class="paramtype">uint&#160;</td>
 348           <td class="paramname"><em>src_stride_x</em>, </td>
 349         </tr>
 350         <tr>
 351           <td class="paramkey"></td>
 352           <td></td>
 353           <td class="paramtype">uint&#160;</td>
 354           <td class="paramname"><em>src_step_x</em>, </td>
 355         </tr>
 356         <tr>
 357           <td class="paramkey"></td>
 358           <td></td>
 359           <td class="paramtype">uint&#160;</td>
 360           <td class="paramname"><em>src_stride_y</em>, </td>
 361         </tr>
 362         <tr>
 363           <td class="paramkey"></td>
 364           <td></td>
 365           <td class="paramtype">uint&#160;</td>
 366           <td class="paramname"><em>src_step_y</em>, </td>
 367         </tr>
 368         <tr>
 369           <td class="paramkey"></td>
 370           <td></td>
 371           <td class="paramtype">uint&#160;</td>
 372           <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
 373         </tr>
 374         <tr>
 375           <td class="paramkey"></td>
 376           <td></td>
 377           <td class="paramtype">__global uchar *&#160;</td>
 378           <td class="paramname"><em>sum_ptr</em>, </td>
 379         </tr>
 380         <tr>
 381           <td class="paramkey"></td>
 382           <td></td>
 383           <td class="paramtype">uint&#160;</td>
 384           <td class="paramname"><em>sum_stride_x</em>, </td>
 385         </tr>
 386         <tr>
 387           <td class="paramkey"></td>
 388           <td></td>
 389           <td class="paramtype">uint&#160;</td>
 390           <td class="paramname"><em>sum_step_x</em>, </td>
 391         </tr>
 392         <tr>
 393           <td class="paramkey"></td>
 394           <td></td>
 395           <td class="paramtype">uint&#160;</td>
 396           <td class="paramname"><em>sum_stride_y</em>, </td>
 397         </tr>
 398         <tr>
 399           <td class="paramkey"></td>
 400           <td></td>
 401           <td class="paramtype">uint&#160;</td>
 402           <td class="paramname"><em>sum_step_y</em>, </td>
 403         </tr>
 404         <tr>
 405           <td class="paramkey"></td>
 406           <td></td>
 407           <td class="paramtype">uint&#160;</td>
 408           <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
 409         </tr>
 410         <tr>
 411           <td class="paramkey"></td>
 412           <td></td>
 413           <td class="paramtype">__global uchar *&#160;</td>
 414           <td class="paramname"><em>dst_ptr</em>, </td>
 415         </tr>
 416         <tr>
 417           <td class="paramkey"></td>
 418           <td></td>
 419           <td class="paramtype">uint&#160;</td>
 420           <td class="paramname"><em>dst_stride_x</em>, </td>
 421         </tr>
 422         <tr>
 423           <td class="paramkey"></td>
 424           <td></td>
 425           <td class="paramtype">uint&#160;</td>
 426           <td class="paramname"><em>dst_step_x</em>, </td>
 427         </tr>
 428         <tr>
 429           <td class="paramkey"></td>
 430           <td></td>
 431           <td class="paramtype">uint&#160;</td>
 432           <td class="paramname"><em>dst_stride_y</em>, </td>
 433         </tr>
 434         <tr>
 435           <td class="paramkey"></td>
 436           <td></td>
 437           <td class="paramtype">uint&#160;</td>
 438           <td class="paramname"><em>dst_step_y</em>, </td>
 439         </tr>
 440         <tr>
 441           <td class="paramkey"></td>
 442           <td></td>
 443           <td class="paramtype">uint&#160;</td>
 444           <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
 445         </tr>
 446         <tr>
 447           <td></td>
 448           <td>)</td>
 449           <td></td><td></td>
 450         </tr>
 451       </table>
 452 </div><div class="memdoc">
 453
 454 <p>Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel. </p>
 455 <dl class="section note"><dt>Note</dt><dd>Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short</dd></dl>
 456 <dl class="params"><dt>Parameters</dt><dd>
 457   <table class="params">
 458     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16, F32 </td></tr>
 459     <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>
 460     <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>
 461     <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>
 462     <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>
 463     <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>
 464     <tr><td class="paramdir">[in]</td><td class="paramname">sum_ptr</td><td>Pointer to the sum values tensor slice. Supported data types: F16, F32 </td></tr>
 465     <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_x</td><td>Stride of the sum values tensor in X dimension (in bytes) </td></tr>
 466     <tr><td class="paramdir">[in]</td><td class="paramname">sum_step_x</td><td>sum_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 467     <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_y</td><td>Stride of the sum values tensor in Y dimension (in bytes) </td></tr>
 468     <tr><td class="paramdir">[in]</td><td class="paramname">sum_step_y</td><td>sum_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 469     <tr><td class="paramdir">[in]</td><td class="paramname">sum_offset_first_element_in_bytes</td><td>The offset of the first element in the sum values tensor </td></tr>
 470     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor slice. Supported data types: F16, F32 </td></tr>
 471     <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>
 472     <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>
 473     <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>
 474     <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>
 475     <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>
 476   </table>
 477   </dd>
 478 </dl>
 479
 480 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00207">207</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 481
 482 <p>References <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00076">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00033">DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00201">offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00032">VEC_DATA_TYPE</a>.</p>
 483 <div class="fragment"><div class="line"><a name="l00211"></a><span class="lineno">  211</span>&#160;{</div><div class="line"><a name="l00212"></a><span class="lineno">  212</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="l00213"></a><span class="lineno">  213</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="l00214"></a><span class="lineno">  214</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> sum = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(sum);</div><div class="line"><a name="l00215"></a><span class="lineno">  215</span>&#160;</div><div class="line"><a name="l00216"></a><span class="lineno">  216</span>&#160;    <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00217"></a><span class="lineno">  217</span>&#160;    <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> sum_val = *((__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;sum, 0, get_global_id(1)));</div><div class="line"><a name="l00218"></a><span class="lineno">  218</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00219"></a><span class="lineno">  219</span>&#160;    data = vload16(0, (__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0));</div><div class="line"><a name="l00220"></a><span class="lineno">  220</span>&#160;    vstore16(data / sum_val, 0, (__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, 0, 0));</div><div class="line"><a name="l00221"></a><span class="lineno">  221</span>&#160;}</div><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>
 484 <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>
 485 <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>
 486 <div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00032">helpers.h:32</a></div></div>
 487 <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>
 488 <div class="ttc" id="softmax__layer_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00033">softmax_layer.cl:33</a></div></div>
 489 </div><!-- fragment -->
 490 </div>
 491 </div>
 492 <a class="anchor" id="aaf2858ffa79555d18154eae8a32db43e"></a>
 493 <div class="memitem">
 494 <div class="memproto">
 495       <table class="memname">
 496         <tr>
 497           <td class="memname">__kernel void softmax_layer_shift_exp_sum </td>
 498           <td>(</td>
 499           <td class="paramtype">__global uchar *&#160;</td>
 500           <td class="paramname"><em>src_ptr</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_stride_x</em>, </td>
 507         </tr>
 508         <tr>
 509           <td class="paramkey"></td>
 510           <td></td>
 511           <td class="paramtype">uint&#160;</td>
 512           <td class="paramname"><em>src_step_x</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>src_stride_y</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>src_step_y</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>src_offset_first_element_in_bytes</em>, </td>
 531         </tr>
 532         <tr>
 533           <td class="paramkey"></td>
 534           <td></td>
 535           <td class="paramtype">__global uchar *&#160;</td>
 536           <td class="paramname"><em>max_ptr</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>max_stride_x</em>, </td>
 543         </tr>
 544         <tr>
 545           <td class="paramkey"></td>
 546           <td></td>
 547           <td class="paramtype">uint&#160;</td>
 548           <td class="paramname"><em>max_step_x</em>, </td>
 549         </tr>
 550         <tr>
 551           <td class="paramkey"></td>
 552           <td></td>
 553           <td class="paramtype">uint&#160;</td>
 554           <td class="paramname"><em>max_stride_y</em>, </td>
 555         </tr>
 556         <tr>
 557           <td class="paramkey"></td>
 558           <td></td>
 559           <td class="paramtype">uint&#160;</td>
 560           <td class="paramname"><em>max_step_y</em>, </td>
 561         </tr>
 562         <tr>
 563           <td class="paramkey"></td>
 564           <td></td>
 565           <td class="paramtype">uint&#160;</td>
 566           <td class="paramname"><em>max_offset_first_element_in_bytes</em>, </td>
 567         </tr>
 568         <tr>
 569           <td class="paramkey"></td>
 570           <td></td>
 571           <td class="paramtype">__global uchar *&#160;</td>
 572           <td class="paramname"><em>dst_ptr</em>, </td>
 573         </tr>
 574         <tr>
 575           <td class="paramkey"></td>
 576           <td></td>
 577           <td class="paramtype">uint&#160;</td>
 578           <td class="paramname"><em>dst_stride_x</em>, </td>
 579         </tr>
 580         <tr>
 581           <td class="paramkey"></td>
 582           <td></td>
 583           <td class="paramtype">uint&#160;</td>
 584           <td class="paramname"><em>dst_step_x</em>, </td>
 585         </tr>
 586         <tr>
 587           <td class="paramkey"></td>
 588           <td></td>
 589           <td class="paramtype">uint&#160;</td>
 590           <td class="paramname"><em>dst_stride_y</em>, </td>
 591         </tr>
 592         <tr>
 593           <td class="paramkey"></td>
 594           <td></td>
 595           <td class="paramtype">uint&#160;</td>
 596           <td class="paramname"><em>dst_step_y</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>dst_offset_first_element_in_bytes</em>, </td>
 603         </tr>
 604         <tr>
 605           <td class="paramkey"></td>
 606           <td></td>
 607           <td class="paramtype">__global uchar *&#160;</td>
 608           <td class="paramname"><em>sum_ptr</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>sum_stride_x</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>sum_step_x</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>sum_stride_y</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>sum_step_y</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>sum_offset_first_element_in_bytes</em>, </td>
 639         </tr>
 640         <tr>
 641           <td class="paramkey"></td>
 642           <td></td>
 643           <td class="paramtype">uint&#160;</td>
 644           <td class="paramname"><em>width</em>&#160;</td>
 645         </tr>
 646         <tr>
 647           <td></td>
 648           <td>)</td>
 649           <td></td><td></td>
 650         </tr>
 651       </table>
 652 </div><div class="memdoc">
 653
 654 <p>Shifts the values of the input tensor by the max calculated in softmax_layer_max kernel, then gets the exponent of each element as sums all elements across each row. </p>
 655 <dl class="section note"><dt>Note</dt><dd>Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short </dd>
 656 <dd>
 657 In case F16 is used -DUSE_HALF must be passed otherwise the kernel will default to used F32. </dd>
 658 <dd>
 659 In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed.</dd></dl>
 660 <dl class="params"><dt>Parameters</dt><dd>
 661   <table class="params">
 662     <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16, F32 </td></tr>
 663     <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>
 664     <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>
 665     <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>
 666     <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>
 667     <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>
 668     <tr><td class="paramdir">[in]</td><td class="paramname">max_ptr</td><td>Pointer to the max values tensor slice. Supported data types: F16, F32 </td></tr>
 669     <tr><td class="paramdir">[in]</td><td class="paramname">max_stride_x</td><td>Stride of the max values tensor in X dimension (in bytes) </td></tr>
 670     <tr><td class="paramdir">[in]</td><td class="paramname">max_step_x</td><td>max_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 671     <tr><td class="paramdir">[in]</td><td class="paramname">max_stride_y</td><td>Stride of the max values tensor in Y dimension (in bytes) </td></tr>
 672     <tr><td class="paramdir">[in]</td><td class="paramname">max_step_y</td><td>max_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 673     <tr><td class="paramdir">[in]</td><td class="paramname">max_offset_first_element_in_bytes</td><td>The offset of the first element in the max values tensor </td></tr>
 674     <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor slice. Supported data types: F16, F32 </td></tr>
 675     <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>
 676     <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>
 677     <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>
 678     <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>
 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 tensor </td></tr>
 680     <tr><td class="paramdir">[out]</td><td class="paramname">sum_ptr</td><td>Pointer to the sum values tensor slice. Supported data types: F16, F32 </td></tr>
 681     <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_x</td><td>Stride of the sum values tensor in X dimension (in bytes) </td></tr>
 682     <tr><td class="paramdir">[in]</td><td class="paramname">sum_step_x</td><td>sum_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
 683     <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_y</td><td>Stride of the sum values tensor in Y dimension (in bytes) </td></tr>
 684     <tr><td class="paramdir">[in]</td><td class="paramname">sum_step_y</td><td>sum_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
 685     <tr><td class="paramdir">[in]</td><td class="paramname">sum_offset_first_element_in_bytes</td><td>The offset of the first element in the sum values tensor </td></tr>
 686     <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
 687   </table>
 688   </dd>
 689 </dl>
 690
 691 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00132">132</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 692
 693 <p>References <a class="el" href="helpers_8h_source.xhtml#l00035">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00073">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00033">DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00037">idx16</a>, <a class="el" href="helpers_8h_source.xhtml#l00201">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00097">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00032">SELECT_DATA_TYPE</a>, and <a class="el" href="helpers_8h_source.xhtml#l00032">VEC_DATA_TYPE</a>.</p>
 694 <div class="fragment"><div class="line"><a name="l00138"></a><span class="lineno">  138</span>&#160;{</div><div class="line"><a name="l00139"></a><span class="lineno">  139</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="l00140"></a><span class="lineno">  140</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="l00141"></a><span class="lineno">  141</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> max = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(max);</div><div class="line"><a name="l00142"></a><span class="lineno">  142</span>&#160;    <a class="code" href="struct_image.xhtml">Image</a> sum = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(sum);</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 max value of 1D logits vector (row)</span></div><div class="line"><a name="l00145"></a><span class="lineno">  145</span>&#160;    <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> max_val = *((__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;max, 0, 0));</div><div class="line"><a name="l00146"></a><span class="lineno">  146</span>&#160;</div><div class="line"><a name="l00147"></a><span class="lineno">  147</span>&#160;    <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00148"></a><span class="lineno">  148</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00149"></a><span class="lineno">  149</span>&#160;    sum1D = 0;</div><div class="line"><a name="l00150"></a><span class="lineno">  150</span>&#160;</div><div class="line"><a name="l00151"></a><span class="lineno">  151</span>&#160;    <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00152"></a><span class="lineno">  152</span>&#160;    const uint width4 = width &gt;&gt; 4;</div><div class="line"><a name="l00153"></a><span class="lineno">  153</span>&#160;    for(uint i = 0; i &lt; width4; i++)</div><div class="line"><a name="l00154"></a><span class="lineno">  154</span>&#160;    {</div><div class="line"><a name="l00155"></a><span class="lineno">  155</span>&#160;        <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;        data = vload16(0, (__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; 4, 0));</div><div class="line"><a name="l00157"></a><span class="lineno">  157</span>&#160;        data = exp(data - max_val);</div><div class="line"><a name="l00158"></a><span class="lineno">  158</span>&#160;        vstore16(data, 0, (__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, i &lt;&lt; 4, 0));</div><div class="line"><a name="l00159"></a><span class="lineno">  159</span>&#160;        sum1D += data;</div><div class="line"><a name="l00160"></a><span class="lineno">  160</span>&#160;    }</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;</div><div class="line"><a name="l00162"></a><span class="lineno">  162</span>&#160;<span class="preprocessor">#if defined NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00163"></a><span class="lineno">  163</span>&#160;    <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00164"></a><span class="lineno">  164</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00165"></a><span class="lineno">  165</span>&#160;    data = vload16(0, (__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width4 &lt;&lt; 4, 0));</div><div class="line"><a name="l00166"></a><span class="lineno">  166</span>&#160;    data = exp(data - max_val);</div><div class="line"><a name="l00167"></a><span class="lineno">  167</span>&#160;    <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 16)</div><div class="line"><a name="l00168"></a><span class="lineno">  168</span>&#160;    widx = <a class="code" href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint16)(width4 &lt;&lt; 4) + <a class="code" href="softmax__layer_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a>) &lt; width, <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 16));</div><div class="line"><a name="l00169"></a><span class="lineno">  169</span>&#160;    data = select(0, data, widx);</div><div class="line"><a name="l00170"></a><span class="lineno">  170</span>&#160;    vstore16(data, 0, (__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, width4 &lt;&lt; 4, 0));</div><div class="line"><a name="l00171"></a><span class="lineno">  171</span>&#160;    sum1D += data;</div><div class="line"><a name="l00172"></a><span class="lineno">  172</span>&#160;<span class="preprocessor">#endif</span></div><div class="line"><a name="l00173"></a><span class="lineno">  173</span>&#160;</div><div class="line"><a name="l00174"></a><span class="lineno">  174</span>&#160;    <span class="comment">// Perform min/max reduction</span></div><div class="line"><a name="l00175"></a><span class="lineno">  175</span>&#160;    sum1D.s01234567 = sum1D.s01234567 + sum1D.s89ABCDEF;</div><div class="line"><a name="l00176"></a><span class="lineno">  176</span>&#160;    sum1D.s0123     = sum1D.s0123 + sum1D.s4567;</div><div class="line"><a name="l00177"></a><span class="lineno">  177</span>&#160;    sum1D.s01       = sum1D.s01 + sum1D.s23;</div><div class="line"><a name="l00178"></a><span class="lineno">  178</span>&#160;    sum1D.s0        = sum1D.s0 + sum1D.s1;</div><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;    <span class="comment">// Calculate and store result</span></div><div class="line"><a name="l00181"></a><span class="lineno">  181</span>&#160;    *((__global <a class="code" href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div><div class="line"><a name="l00182"></a><span class="lineno">  182</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>
 695 <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>
 696 <div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00032">helpers.h:32</a></div></div>
 697 <div class="ttc" id="softmax__layer_8cl_xhtml_af5b2e33e3c5fcaab3a213f26c2300170"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a></div><div class="ttdeci">#define SELECT_DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00032">softmax_layer.cl:32</a></div></div>
 698 <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>
 699 <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>
 700 <div class="ttc" id="softmax__layer_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="softmax__layer_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00033">softmax_layer.cl:33</a></div></div>
 701 <div class="ttc" id="softmax__layer_8cl_xhtml_a0712735973f172ac9efc7d48a31e47ad"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a0712735973f172ac9efc7d48a31e47ad">idx16</a></div><div class="ttdeci">__constant uint16 idx16</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00037">softmax_layer.cl:37</a></div></div>
 702 <div class="ttc" id="helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00035">helpers.h:35</a></div></div>
 703 </div><!-- fragment -->
 704 </div>
 705 </div>
 706 <h2 class="groupheader">Variable Documentation</h2>
 707 <a class="anchor" id="a0712735973f172ac9efc7d48a31e47ad"></a>
 708 <div class="memitem">
 709 <div class="memproto">
 710       <table class="memname">
 711         <tr>
 712           <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
 713         </tr>
 714       </table>
 715 </div><div class="memdoc">
 716
 717 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00037">37</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 718
 719 <p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00059">softmax_layer_max()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00132">softmax_layer_shift_exp_sum()</a>.</p>
 720
 721 </div>
 722 </div>
 723 <a class="anchor" id="a5505428916dfe70b40b32686c57e8436"></a>
 724 <div class="memitem">
 725 <div class="memproto">
 726       <table class="memname">
 727         <tr>
 728           <td class="memname">__constant float16 type_min = ( float16 )( FLT_MIN )</td>
 729         </tr>
 730       </table>
 731 </div><div class="memdoc">
 732
 733 <p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00036">36</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
 734
 735 <p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00059">softmax_layer_max()</a>.</p>
 736
 737 </div>
 738 </div>
 739 </div><!-- contents -->
 740 </div><!-- doc-content -->
 741 <!-- start footer part -->
 742 <div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
 743   <ul>
 744     <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="softmax__layer_8cl.xhtml">softmax_layer.cl</a></li>
 745     <li class="footer">Generated on Wed Apr 12 2017 14:26:05 for ARM Compute Library by
 746     <a href="http://www.doxygen.org/index.html">
 747     <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
 748   </ul>
 749 </div>
 750 </body>
 751 </html>