Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 8938bd3f40ea62ff56d6ed4e2db0a8aee34dd64a / . / documentation / softmax__layer_8cl.xhtml
blob: d90043cb00c27a01d581379cecbe6d6d79c4524f [file] [log] [blame]
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]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"/>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]7<meta name="generator" content="Doxygen 1.8.6"/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]8<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]9<title>Compute Library: src/core/CL/cl_kernels/softmax_layer.cl File Reference</title>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]15<script type="text/javascript" src="navtree.js"></script>
16<script type="text/javascript">
17 $(document).ready(initResizable);
18 $(window).load(resizeHeight);
19</script>
20<link href="search/search.css" rel="stylesheet" type="text/css"/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]21<script type="text/javascript" src="search/search.js"></script>
22<script type="text/javascript">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]23 $(document).ready(function() { searchBox.OnSelectItem(0); });
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]24</script>
25<script type="text/x-mathjax-config">
26 MathJax.Hub.Config({
27 extensions: ["tex2jax.js"],
28 jax: ["input/TeX","output/HTML-CSS"],
29});
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]30</script><script src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]31<link href="doxygen.css" rel="stylesheet" type="text/css" />
32</head>
33<body>
34<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
35<div id="titlearea">
36<table cellspacing="0" cellpadding="0">
37 <tbody>
38 <tr style="height: 56px;">
39 <td style="padding-left: 0.5em;">
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]40 <div id="projectname">Compute Library
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]41 &#160;<span id="projectnumber">17.09</span>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]42 </div>
43 </td>
44 </tr>
45 </tbody>
46</table>
47</div>
48<!-- end header part -->
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]49<!-- Generated by Doxygen 1.8.6 -->
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]50<script type="text/javascript">
51var searchBox = new SearchBox("searchBox", "search",false,'Search');
52</script>
53 <div id="navrow1" class="tabs">
54 <ul class="tablist">
55 <li><a href="index.xhtml"><span>Main&#160;Page</span></a></li>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]56 <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]57 <li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
58 <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
59 <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
60 <li>
61 <div id="MSearchBox" class="MSearchBoxInactive">
62 <span class="left">
63 <img id="MSearchSelect" src="search/mag_sel.png"
64 onmouseover="return searchBox.OnSearchSelectShow()"
65 onmouseout="return searchBox.OnSearchSelectHide()"
66 alt=""/>
67 <input type="text" id="MSearchField" value="Search" accesskey="S"
68 onfocus="searchBox.OnSearchFieldFocus(true)"
69 onblur="searchBox.OnSearchFieldFocus(false)"
70 onkeyup="searchBox.OnSearchFieldChange(event)"/>
71 </span><span class="right">
72 <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
73 </span>
74 </div>
75 </li>
76 </ul>
77 </div>
78 <div id="navrow2" class="tabs2">
79 <ul class="tablist">
80 <li><a href="files.xhtml"><span>File&#160;List</span></a></li>
81 <li><a href="globals.xhtml"><span>Globals</span></a></li>
82 </ul>
83 </div>
84</div><!-- top -->
85<div id="side-nav" class="ui-resizable side-nav-resizable">
86 <div id="nav-tree">
87 <div id="nav-tree-contents">
88 <div id="nav-sync" class="sync"></div>
89 </div>
90 </div>
91 <div id="splitbar" style="-moz-user-select:none;"
92 class="ui-resizable-handle">
93 </div>
94</div>
95<script type="text/javascript">
96$(document).ready(function(){initNavTree('softmax__layer_8cl.xhtml','');});
97</script>
98<div id="doc-content">
99<!-- window showing the filter options -->
100<div id="MSearchSelectWindow"
101 onmouseover="return searchBox.OnSearchSelectShow()"
102 onmouseout="return searchBox.OnSearchSelectHide()"
103 onkeydown="return searchBox.OnSearchSelectKey(event)">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]104<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Data Structures</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Friends</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(10)"><span class="SelectionMark">&#160;</span>Macros</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(11)"><span class="SelectionMark">&#160;</span>Pages</a></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]105
106<!-- iframe showing the search results (closed by default) -->
107<div id="MSearchResultsWindow">
108<iframe src="javascript:void(0)" frameborder="0"
109 name="MSearchResults" id="MSearchResults">
110</iframe>
111</div>
112
113<div class="header">
114 <div class="summary">
115<a href="#define-members">Macros</a> &#124;
116<a href="#func-members">Functions</a> &#124;
117<a href="#var-members">Variables</a> </div>
118 <div class="headertitle">
119<div class="title">softmax_layer.cl File Reference</div> </div>
120</div><!--header-->
121<div class="contents">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]122<div class="textblock"><code>#include &quot;<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]123</div>
124<p><a href="softmax__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
125<table class="memberdecls">
126<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
127Macros</h2></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]128<tr class="memitem:abaa48ad818c44e415fd3f9dd0f27bf01"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(x, y, type, <a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>)&#160;&#160;&#160;max((x), (y))</td></tr>
129<tr class="separator:abaa48ad818c44e415fd3f9dd0f27bf01"><td class="memSeparator" colspan="2">&#160;</td></tr>
130<tr class="memitem:a44206a4e5783c7aabacec88aad878c88"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(x, y, type, <a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>)&#160;&#160;&#160;((x) + (y))</td></tr>
131<tr class="separator:a44206a4e5783c7aabacec88aad878c88"><td class="memSeparator" colspan="2">&#160;</td></tr>
132<tr class="memitem:ac3af2d18008cbbf7247ae48fcd6e0c4e"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(x, y, type, <a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>)&#160;&#160;&#160;((x) - (y))</td></tr>
133<tr class="separator:ac3af2d18008cbbf7247ae48fcd6e0c4e"><td class="memSeparator" colspan="2">&#160;</td></tr>
134<tr class="memitem:a8cde99b1ce0f3c1dacd49261b0cf03d8"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a8cde99b1ce0f3c1dacd49261b0cf03d8">DIV_OP</a>(x, y, type, <a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>)&#160;&#160;&#160;((x) / (y))</td></tr>
135<tr class="separator:a8cde99b1ce0f3c1dacd49261b0cf03d8"><td class="memSeparator" colspan="2">&#160;</td></tr>
136<tr class="memitem:a93cf800667317d96574477b9f0a75234"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(x, type, <a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>)&#160;&#160;&#160;exp((x))</td></tr>
137<tr class="separator:a93cf800667317d96574477b9f0a75234"><td class="memSeparator" colspan="2">&#160;</td></tr>
138<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_MAX</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]139<tr class="separator:a80b22c555ddadb47cc6ca338a9c49126"><td class="memSeparator" colspan="2">&#160;</td></tr>
140<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>
141<tr class="separator:af5b2e33e3c5fcaab3a213f26c2300170"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]142</table><table class="memberdecls">
143<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
144Functions</h2></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]145<tr class="memitem:a9e86f72ca0d249cb88c36bd50620da8a"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a9e86f72ca0d249cb88c36bd50620da8a">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_stride_z, uint src_step_z, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes, uint width)</td></tr>
146<tr class="memdesc:a9e86f72ca0d249cb88c36bd50620da8a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Identifies the maximum value across the 1st dimension. <a href="#a9e86f72ca0d249cb88c36bd50620da8a">More...</a><br/></td></tr>
147<tr class="separator:a9e86f72ca0d249cb88c36bd50620da8a"><td class="memSeparator" colspan="2">&#160;</td></tr>
148<tr class="memitem:aa4412bb319c67deaeab07a481d93dfdb"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#aa4412bb319c67deaeab07a481d93dfdb">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_stride_z, uint src_step_z, 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_stride_z, uint max_step_z, 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_stride_z, uint dst_step_z, 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_stride_z, uint sum_step_z, uint sum_offset_first_element_in_bytes, uint width)</td></tr>
149<tr class="memdesc:aa4412bb319c67deaeab07a481d93dfdb"><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="#aa4412bb319c67deaeab07a481d93dfdb">More...</a><br/></td></tr>
150<tr class="separator:aa4412bb319c67deaeab07a481d93dfdb"><td class="memSeparator" colspan="2">&#160;</td></tr>
151<tr class="memitem:ac4247ac0991e85965b7ded764e78f12c"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#ac4247ac0991e85965b7ded764e78f12c">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_stride_z, uint src_step_z, 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_stride_z, uint sum_step_z, 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_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes)</td></tr>
152<tr class="memdesc:ac4247ac0991e85965b7ded764e78f12c"><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="#ac4247ac0991e85965b7ded764e78f12c">More...</a><br/></td></tr>
153<tr class="separator:ac4247ac0991e85965b7ded764e78f12c"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]154</table><table class="memberdecls">
155<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="var-members"></a>
156Variables</h2></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]157<tr class="memitem:a538b4b63f40e7b12891774e03a4f0dec"><td class="memItemLeft" align="right" valign="top">__constant DATA_TYPE16&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a> = ( DATA_TYPE16 )( -FLT_MAX )</td></tr>
158<tr class="separator:a538b4b63f40e7b12891774e03a4f0dec"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]159<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>
160<tr class="separator:a0712735973f172ac9efc7d48a31e47ad"><td class="memSeparator" colspan="2">&#160;</td></tr>
161</table>
162<h2 class="groupheader">Macro Definition Documentation</h2>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]163<a class="anchor" id="a44206a4e5783c7aabacec88aad878c88"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]164<div class="memitem">
165<div class="memproto">
166 <table class="memname">
167 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]168 <td class="memname">#define ADD_OP</td>
169 <td>(</td>
170 <td class="paramtype">&#160;</td>
171 <td class="paramname">x, </td>
172 </tr>
173 <tr>
174 <td class="paramkey"></td>
175 <td></td>
176 <td class="paramtype">&#160;</td>
177 <td class="paramname">y, </td>
178 </tr>
179 <tr>
180 <td class="paramkey"></td>
181 <td></td>
182 <td class="paramtype">&#160;</td>
183 <td class="paramname">type, </td>
184 </tr>
185 <tr>
186 <td class="paramkey"></td>
187 <td></td>
188 <td class="paramtype">&#160;</td>
189 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
190 </tr>
191 <tr>
192 <td></td>
193 <td>)</td>
194 <td></td><td>&#160;&#160;&#160;((x) + (y))</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]195 </tr>
196 </table>
197</div><div class="memdoc">
198
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]199<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">43</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]200
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]201<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00166">softmax_layer_shift_exp_sum()</a>.</p>
202
203</div>
204</div>
205<a class="anchor" id="a8cde99b1ce0f3c1dacd49261b0cf03d8"></a>
206<div class="memitem">
207<div class="memproto">
208 <table class="memname">
209 <tr>
210 <td class="memname">#define DIV_OP</td>
211 <td>(</td>
212 <td class="paramtype">&#160;</td>
213 <td class="paramname">x, </td>
214 </tr>
215 <tr>
216 <td class="paramkey"></td>
217 <td></td>
218 <td class="paramtype">&#160;</td>
219 <td class="paramname">y, </td>
220 </tr>
221 <tr>
222 <td class="paramkey"></td>
223 <td></td>
224 <td class="paramtype">&#160;</td>
225 <td class="paramname">type, </td>
226 </tr>
227 <tr>
228 <td class="paramkey"></td>
229 <td></td>
230 <td class="paramtype">&#160;</td>
231 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
232 </tr>
233 <tr>
234 <td></td>
235 <td>)</td>
236 <td></td><td>&#160;&#160;&#160;((x) / (y))</td>
237 </tr>
238 </table>
239</div><div class="memdoc">
240
241<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">45</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
242
243<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00250">softmax_layer_norm()</a>.</p>
244
245</div>
246</div>
247<a class="anchor" id="a93cf800667317d96574477b9f0a75234"></a>
248<div class="memitem">
249<div class="memproto">
250 <table class="memname">
251 <tr>
252 <td class="memname">#define EXP_OP</td>
253 <td>(</td>
254 <td class="paramtype">&#160;</td>
255 <td class="paramname">x, </td>
256 </tr>
257 <tr>
258 <td class="paramkey"></td>
259 <td></td>
260 <td class="paramtype">&#160;</td>
261 <td class="paramname">type, </td>
262 </tr>
263 <tr>
264 <td class="paramkey"></td>
265 <td></td>
266 <td class="paramtype">&#160;</td>
267 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
268 </tr>
269 <tr>
270 <td></td>
271 <td>)</td>
272 <td></td><td>&#160;&#160;&#160;exp((x))</td>
273 </tr>
274 </table>
275</div><div class="memdoc">
276
277<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00046">46</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
278
279<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00166">softmax_layer_shift_exp_sum()</a>.</p>
280
281</div>
282</div>
283<a class="anchor" id="abaa48ad818c44e415fd3f9dd0f27bf01"></a>
284<div class="memitem">
285<div class="memproto">
286 <table class="memname">
287 <tr>
288 <td class="memname">#define MAX_OP</td>
289 <td>(</td>
290 <td class="paramtype">&#160;</td>
291 <td class="paramname">x, </td>
292 </tr>
293 <tr>
294 <td class="paramkey"></td>
295 <td></td>
296 <td class="paramtype">&#160;</td>
297 <td class="paramname">y, </td>
298 </tr>
299 <tr>
300 <td class="paramkey"></td>
301 <td></td>
302 <td class="paramtype">&#160;</td>
303 <td class="paramname">type, </td>
304 </tr>
305 <tr>
306 <td class="paramkey"></td>
307 <td></td>
308 <td class="paramtype">&#160;</td>
309 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
310 </tr>
311 <tr>
312 <td></td>
313 <td>)</td>
314 <td></td><td>&#160;&#160;&#160;max((x), (y))</td>
315 </tr>
316 </table>
317</div><div class="memdoc">
318
319<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00042">42</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
320
321<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00085">softmax_layer_max()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]322
323</div>
324</div>
325<a class="anchor" id="a80b22c555ddadb47cc6ca338a9c49126"></a>
326<div class="memitem">
327<div class="memproto">
328 <table class="memname">
329 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]330 <td class="memname">#define MINVAL&#160;&#160;&#160;-FLT_MAX</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]331 </tr>
332 </table>
333</div><div class="memdoc">
334
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]335<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00052">52</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]336
337</div>
338</div>
339<a class="anchor" id="af5b2e33e3c5fcaab3a213f26c2300170"></a>
340<div class="memitem">
341<div class="memproto">
342 <table class="memname">
343 <tr>
344 <td class="memname">#define SELECT_DATA_TYPE&#160;&#160;&#160;int</td>
345 </tr>
346 </table>
347</div><div class="memdoc">
348
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]349<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00053">53</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]350
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]351<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00085">softmax_layer_max()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00166">softmax_layer_shift_exp_sum()</a>.</p>
352
353</div>
354</div>
355<a class="anchor" id="ac3af2d18008cbbf7247ae48fcd6e0c4e"></a>
356<div class="memitem">
357<div class="memproto">
358 <table class="memname">
359 <tr>
360 <td class="memname">#define SUB_OP</td>
361 <td>(</td>
362 <td class="paramtype">&#160;</td>
363 <td class="paramname">x, </td>
364 </tr>
365 <tr>
366 <td class="paramkey"></td>
367 <td></td>
368 <td class="paramtype">&#160;</td>
369 <td class="paramname">y, </td>
370 </tr>
371 <tr>
372 <td class="paramkey"></td>
373 <td></td>
374 <td class="paramtype">&#160;</td>
375 <td class="paramname">type, </td>
376 </tr>
377 <tr>
378 <td class="paramkey"></td>
379 <td></td>
380 <td class="paramtype">&#160;</td>
381 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
382 </tr>
383 <tr>
384 <td></td>
385 <td>)</td>
386 <td></td><td>&#160;&#160;&#160;((x) - (y))</td>
387 </tr>
388 </table>
389</div><div class="memdoc">
390
391<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">44</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
392
393<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00166">softmax_layer_shift_exp_sum()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]394
395</div>
396</div>
397<h2 class="groupheader">Function Documentation</h2>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]398<a class="anchor" id="a9e86f72ca0d249cb88c36bd50620da8a"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]399<div class="memitem">
400<div class="memproto">
401 <table class="memname">
402 <tr>
403 <td class="memname">__kernel void softmax_layer_max </td>
404 <td>(</td>
405 <td class="paramtype">__global uchar *&#160;</td>
406 <td class="paramname"><em>src_ptr</em>, </td>
407 </tr>
408 <tr>
409 <td class="paramkey"></td>
410 <td></td>
411 <td class="paramtype">uint&#160;</td>
412 <td class="paramname"><em>src_stride_x</em>, </td>
413 </tr>
414 <tr>
415 <td class="paramkey"></td>
416 <td></td>
417 <td class="paramtype">uint&#160;</td>
418 <td class="paramname"><em>src_step_x</em>, </td>
419 </tr>
420 <tr>
421 <td class="paramkey"></td>
422 <td></td>
423 <td class="paramtype">uint&#160;</td>
424 <td class="paramname"><em>src_stride_y</em>, </td>
425 </tr>
426 <tr>
427 <td class="paramkey"></td>
428 <td></td>
429 <td class="paramtype">uint&#160;</td>
430 <td class="paramname"><em>src_step_y</em>, </td>
431 </tr>
432 <tr>
433 <td class="paramkey"></td>
434 <td></td>
435 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]436 <td class="paramname"><em>src_stride_z</em>, </td>
437 </tr>
438 <tr>
439 <td class="paramkey"></td>
440 <td></td>
441 <td class="paramtype">uint&#160;</td>
442 <td class="paramname"><em>src_step_z</em>, </td>
443 </tr>
444 <tr>
445 <td class="paramkey"></td>
446 <td></td>
447 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]448 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
449 </tr>
450 <tr>
451 <td class="paramkey"></td>
452 <td></td>
453 <td class="paramtype">__global uchar *&#160;</td>
454 <td class="paramname"><em>dst_ptr</em>, </td>
455 </tr>
456 <tr>
457 <td class="paramkey"></td>
458 <td></td>
459 <td class="paramtype">uint&#160;</td>
460 <td class="paramname"><em>dst_stride_x</em>, </td>
461 </tr>
462 <tr>
463 <td class="paramkey"></td>
464 <td></td>
465 <td class="paramtype">uint&#160;</td>
466 <td class="paramname"><em>dst_step_x</em>, </td>
467 </tr>
468 <tr>
469 <td class="paramkey"></td>
470 <td></td>
471 <td class="paramtype">uint&#160;</td>
472 <td class="paramname"><em>dst_stride_y</em>, </td>
473 </tr>
474 <tr>
475 <td class="paramkey"></td>
476 <td></td>
477 <td class="paramtype">uint&#160;</td>
478 <td class="paramname"><em>dst_step_y</em>, </td>
479 </tr>
480 <tr>
481 <td class="paramkey"></td>
482 <td></td>
483 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]484 <td class="paramname"><em>dst_stride_z</em>, </td>
485 </tr>
486 <tr>
487 <td class="paramkey"></td>
488 <td></td>
489 <td class="paramtype">uint&#160;</td>
490 <td class="paramname"><em>dst_step_z</em>, </td>
491 </tr>
492 <tr>
493 <td class="paramkey"></td>
494 <td></td>
495 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]496 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
497 </tr>
498 <tr>
499 <td class="paramkey"></td>
500 <td></td>
501 <td class="paramtype">uint&#160;</td>
502 <td class="paramname"><em>width</em>&#160;</td>
503 </tr>
504 <tr>
505 <td></td>
506 <td>)</td>
507 <td></td><td></td>
508 </tr>
509 </table>
510</div><div class="memdoc">
511
512<p>Identifies the maximum value across the 1st dimension. </p>
513<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>
514<dd>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]515Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]516<dd>
517In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed.</dd></dl>
518<dl class="params"><dt>Parameters</dt><dd>
519 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]520 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: QS8/QS16/F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]521 <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>
522 <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>
523 <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>
524 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]525 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
526 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]527 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]528 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]529 <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>
530 <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>
531 <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>
532 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]533 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
534 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]535 <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>
536 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
537 </table>
538 </dd>
539</dl>
540
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]541<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00085">85</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]542
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]543<p>References <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00105">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00059">idx16</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00042">MAX_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00292">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00133">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00053">SELECT_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00058">type_min</a>, and <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>.</p>
544<div class="fragment"><div class="line"><a name="l00089"></a><span class="lineno"> 89</span>&#160;{</div>
545<div class="line"><a name="l00090"></a><span class="lineno"> 90</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div>
546<div class="line"><a name="l00091"></a><span class="lineno"> 91</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> dst = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div>
547<div class="line"><a name="l00092"></a><span class="lineno"> 92</span>&#160;</div>
548<div class="line"><a name="l00093"></a><span class="lineno"> 93</span>&#160; <span class="comment">// Initialize local maximum</span></div>
549<div class="line"><a name="l00094"></a><span class="lineno"> 94</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div>
550<div class="line"><a name="l00095"></a><span class="lineno"> 95</span>&#160; max_val = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16))<a class="code" href="softmax__layer_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a>;</div>
551<div class="line"><a name="l00096"></a><span class="lineno"> 96</span>&#160;</div>
552<div class="line"><a name="l00097"></a><span class="lineno"> 97</span>&#160; <span class="comment">// Calculate max of row</span></div>
553<div class="line"><a name="l00098"></a><span class="lineno"> 98</span>&#160; const uint width4 = width &gt;&gt; 4;</div>
554<div class="line"><a name="l00099"></a><span class="lineno"> 99</span>&#160; for(uint i = 0; i &lt; width4; i++)</div>
555<div class="line"><a name="l00100"></a><span class="lineno"> 100</span>&#160; {</div>
556<div class="line"><a name="l00101"></a><span class="lineno"> 101</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div>
557<div class="line"><a name="l00102"></a><span class="lineno"> 102</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>
558<div class="line"><a name="l00103"></a><span class="lineno"> 103</span>&#160; max_val = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data, max_val, DATA_TYPE, 16);</div>
559<div class="line"><a name="l00104"></a><span class="lineno"> 104</span>&#160; }</div>
560<div class="line"><a name="l00105"></a><span class="lineno"> 105</span>&#160;</div>
561<div class="line"><a name="l00106"></a><span class="lineno"> 106</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div>
562<div class="line"><a name="l00107"></a><span class="lineno"> 107</span>&#160;<span class="preprocessor"></span> <span class="comment">// Handle non multiple of 16</span></div>
563<div class="line"><a name="l00108"></a><span class="lineno"> 108</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div>
564<div class="line"><a name="l00109"></a><span class="lineno"> 109</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>
565<div class="line"><a name="l00110"></a><span class="lineno"> 110</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 16)</div>
566<div class="line"><a name="l00111"></a><span class="lineno"> 111</span>&#160; widx = <a class="code" href="fixed__point_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="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 16));</div>
567<div class="line"><a name="l00112"></a><span class="lineno"> 112</span>&#160; max_val = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val, select(type_min, data, widx), DATA_TYPE, 16);</div>
568<div class="line"><a name="l00113"></a><span class="lineno"> 113</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_16 */</span><span class="preprocessor"></span></div>
569<div class="line"><a name="l00114"></a><span class="lineno"> 114</span>&#160;<span class="preprocessor"></span></div>
570<div class="line"><a name="l00115"></a><span class="lineno"> 115</span>&#160; <span class="comment">// Perform max reduction</span></div>
571<div class="line"><a name="l00116"></a><span class="lineno"> 116</span>&#160; max_val.s01234567 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val.s01234567, max_val.s89ABCDEF, DATA_TYPE, 8);</div>
572<div class="line"><a name="l00117"></a><span class="lineno"> 117</span>&#160; max_val.s0123 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val.s0123, max_val.s4567, DATA_TYPE, 4);</div>
573<div class="line"><a name="l00118"></a><span class="lineno"> 118</span>&#160; max_val.s01 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val.s01, max_val.s23, DATA_TYPE, 2);</div>
574<div class="line"><a name="l00119"></a><span class="lineno"> 119</span>&#160; max_val.s0 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val.s0, max_val.s1, DATA_TYPE, 1);</div>
575<div class="line"><a name="l00120"></a><span class="lineno"> 120</span>&#160;</div>
576<div class="line"><a name="l00121"></a><span class="lineno"> 121</span>&#160; <span class="comment">// Store result</span></div>
577<div class="line"><a name="l00122"></a><span class="lineno"> 122</span>&#160; *((__global DATA_TYPE *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = max_val.s0;</div>
578<div class="line"><a name="l00123"></a><span class="lineno"> 123</span>&#160;}</div>
579<div class="ttc" id="fixed__point_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00098">fixed_point.h:98</a></div></div>
580<div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
581<div class="ttc" id="softmax__layer_8cl_xhtml_abaa48ad818c44e415fd3f9dd0f27bf01"><div class="ttname"><a href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a></div><div class="ttdeci">#define MAX_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00042">softmax_layer.cl:42</a></div></div>
582<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#l00292">helpers.h:292</a></div></div>
583<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#l00053">softmax_layer.cl:53</a></div></div>
584<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#l00131">helpers.h:131</a></div></div>
585<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#l00133">helpers.h:133</a></div></div>
586<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_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="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
587<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#l00059">softmax_layer.cl:59</a></div></div>
588<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00105">helpers.h:105</a></div></div>
589<div class="ttc" id="softmax__layer_8cl_xhtml_a538b4b63f40e7b12891774e03a4f0dec"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a></div><div class="ttdeci">__constant DATA_TYPE16 type_min</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00058">softmax_layer.cl:58</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]590</div><!-- fragment -->
591</div>
592</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]593<a class="anchor" id="ac4247ac0991e85965b7ded764e78f12c"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]594<div class="memitem">
595<div class="memproto">
596 <table class="memname">
597 <tr>
598 <td class="memname">__kernel void softmax_layer_norm </td>
599 <td>(</td>
600 <td class="paramtype">__global uchar *&#160;</td>
601 <td class="paramname"><em>src_ptr</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_x</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_x</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_stride_y</em>, </td>
620 </tr>
621 <tr>
622 <td class="paramkey"></td>
623 <td></td>
624 <td class="paramtype">uint&#160;</td>
625 <td class="paramname"><em>src_step_y</em>, </td>
626 </tr>
627 <tr>
628 <td class="paramkey"></td>
629 <td></td>
630 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]631 <td class="paramname"><em>src_stride_z</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>src_step_z</em>, </td>
638 </tr>
639 <tr>
640 <td class="paramkey"></td>
641 <td></td>
642 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]643 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
644 </tr>
645 <tr>
646 <td class="paramkey"></td>
647 <td></td>
648 <td class="paramtype">__global uchar *&#160;</td>
649 <td class="paramname"><em>sum_ptr</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>sum_stride_x</em>, </td>
656 </tr>
657 <tr>
658 <td class="paramkey"></td>
659 <td></td>
660 <td class="paramtype">uint&#160;</td>
661 <td class="paramname"><em>sum_step_x</em>, </td>
662 </tr>
663 <tr>
664 <td class="paramkey"></td>
665 <td></td>
666 <td class="paramtype">uint&#160;</td>
667 <td class="paramname"><em>sum_stride_y</em>, </td>
668 </tr>
669 <tr>
670 <td class="paramkey"></td>
671 <td></td>
672 <td class="paramtype">uint&#160;</td>
673 <td class="paramname"><em>sum_step_y</em>, </td>
674 </tr>
675 <tr>
676 <td class="paramkey"></td>
677 <td></td>
678 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]679 <td class="paramname"><em>sum_stride_z</em>, </td>
680 </tr>
681 <tr>
682 <td class="paramkey"></td>
683 <td></td>
684 <td class="paramtype">uint&#160;</td>
685 <td class="paramname"><em>sum_step_z</em>, </td>
686 </tr>
687 <tr>
688 <td class="paramkey"></td>
689 <td></td>
690 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]691 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
692 </tr>
693 <tr>
694 <td class="paramkey"></td>
695 <td></td>
696 <td class="paramtype">__global uchar *&#160;</td>
697 <td class="paramname"><em>dst_ptr</em>, </td>
698 </tr>
699 <tr>
700 <td class="paramkey"></td>
701 <td></td>
702 <td class="paramtype">uint&#160;</td>
703 <td class="paramname"><em>dst_stride_x</em>, </td>
704 </tr>
705 <tr>
706 <td class="paramkey"></td>
707 <td></td>
708 <td class="paramtype">uint&#160;</td>
709 <td class="paramname"><em>dst_step_x</em>, </td>
710 </tr>
711 <tr>
712 <td class="paramkey"></td>
713 <td></td>
714 <td class="paramtype">uint&#160;</td>
715 <td class="paramname"><em>dst_stride_y</em>, </td>
716 </tr>
717 <tr>
718 <td class="paramkey"></td>
719 <td></td>
720 <td class="paramtype">uint&#160;</td>
721 <td class="paramname"><em>dst_step_y</em>, </td>
722 </tr>
723 <tr>
724 <td class="paramkey"></td>
725 <td></td>
726 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]727 <td class="paramname"><em>dst_stride_z</em>, </td>
728 </tr>
729 <tr>
730 <td class="paramkey"></td>
731 <td></td>
732 <td class="paramtype">uint&#160;</td>
733 <td class="paramname"><em>dst_step_z</em>, </td>
734 </tr>
735 <tr>
736 <td class="paramkey"></td>
737 <td></td>
738 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]739 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
740 </tr>
741 <tr>
742 <td></td>
743 <td>)</td>
744 <td></td><td></td>
745 </tr>
746 </table>
747</div><div class="memdoc">
748
749<p>Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel. </p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]750<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>
751<dd>
752Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4</dd></dl>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]753<dl class="params"><dt>Parameters</dt><dd>
754 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]755 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: QS8/QS16/F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]756 <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>
757 <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>
758 <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>
759 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]760 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
761 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]762 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]763 <tr><td class="paramdir">[in]</td><td class="paramname">sum_ptr</td><td>Pointer to the sum values tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]764 <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>
765 <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>
766 <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>
767 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]768 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_z</td><td>Stride of the sum values tensor in Z dimension (in bytes) </td></tr>
769 <tr><td class="paramdir">[in]</td><td class="paramname">sum_step_z</td><td>sum_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]770 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]771 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]772 <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>
773 <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>
774 <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>
775 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]776 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
777 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]778 <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>
779 </table>
780 </dd>
781</dl>
782
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]783<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00250">250</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]784
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]785<p>References <a class="el" href="helpers_8h_source.xhtml#l00105">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00102">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">DIV_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00292">offset()</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, and <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>.</p>
786<div class="fragment"><div class="line"><a name="l00254"></a><span class="lineno"> 254</span>&#160;{</div>
787<div class="line"><a name="l00255"></a><span class="lineno"> 255</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div>
788<div class="line"><a name="l00256"></a><span class="lineno"> 256</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> dst = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div>
789<div class="line"><a name="l00257"></a><span class="lineno"> 257</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a> = <a class="code" href="helpers_8h.xhtml#a071aa45af973feac43b14f62e54a6fce">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a>(sum);</div>
790<div class="line"><a name="l00258"></a><span class="lineno"> 258</span>&#160;</div>
791<div class="line"><a name="l00259"></a><span class="lineno"> 259</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div>
792<div class="line"><a name="l00260"></a><span class="lineno"> 260</span>&#160; <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> sum_val = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;sum, 0, get_global_id(1)));</div>
793<div class="line"><a name="l00261"></a><span class="lineno"> 261</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div>
794<div class="line"><a name="l00262"></a><span class="lineno"> 262</span>&#160; data = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0));</div>
795<div class="line"><a name="l00263"></a><span class="lineno"> 263</span>&#160; vstore16(<a class="code" href="softmax__layer_8cl.xhtml#a8cde99b1ce0f3c1dacd49261b0cf03d8">DIV_OP</a>(data, sum_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16), 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, 0, 0));</div>
796<div class="line"><a name="l00264"></a><span class="lineno"> 264</span>&#160;}</div>
797<div class="ttc" id="helpers_8h_xhtml_a071aa45af973feac43b14f62e54a6fce"><div class="ttname"><a href="helpers_8h.xhtml#a071aa45af973feac43b14f62e54a6fce">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00102">helpers.h:102</a></div></div>
798<div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
799<div class="ttc" id="reduction__operation_8cl_xhtml_ab0df00f5333da51860deb93deb44a782"><div class="ttname"><a href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a></div><div class="ttdeci">DATA_TYPE sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate sum of a vector. </div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00052">reduction_operation.cl:52</a></div></div>
800<div class="ttc" id="softmax__layer_8cl_xhtml_a8cde99b1ce0f3c1dacd49261b0cf03d8"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a8cde99b1ce0f3c1dacd49261b0cf03d8">DIV_OP</a></div><div class="ttdeci">#define DIV_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00045">softmax_layer.cl:45</a></div></div>
801<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#l00292">helpers.h:292</a></div></div>
802<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#l00131">helpers.h:131</a></div></div>
803<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_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="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
804<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00105">helpers.h:105</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]805</div><!-- fragment -->
806</div>
807</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]808<a class="anchor" id="aa4412bb319c67deaeab07a481d93dfdb"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]809<div class="memitem">
810<div class="memproto">
811 <table class="memname">
812 <tr>
813 <td class="memname">__kernel void softmax_layer_shift_exp_sum </td>
814 <td>(</td>
815 <td class="paramtype">__global uchar *&#160;</td>
816 <td class="paramname"><em>src_ptr</em>, </td>
817 </tr>
818 <tr>
819 <td class="paramkey"></td>
820 <td></td>
821 <td class="paramtype">uint&#160;</td>
822 <td class="paramname"><em>src_stride_x</em>, </td>
823 </tr>
824 <tr>
825 <td class="paramkey"></td>
826 <td></td>
827 <td class="paramtype">uint&#160;</td>
828 <td class="paramname"><em>src_step_x</em>, </td>
829 </tr>
830 <tr>
831 <td class="paramkey"></td>
832 <td></td>
833 <td class="paramtype">uint&#160;</td>
834 <td class="paramname"><em>src_stride_y</em>, </td>
835 </tr>
836 <tr>
837 <td class="paramkey"></td>
838 <td></td>
839 <td class="paramtype">uint&#160;</td>
840 <td class="paramname"><em>src_step_y</em>, </td>
841 </tr>
842 <tr>
843 <td class="paramkey"></td>
844 <td></td>
845 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]846 <td class="paramname"><em>src_stride_z</em>, </td>
847 </tr>
848 <tr>
849 <td class="paramkey"></td>
850 <td></td>
851 <td class="paramtype">uint&#160;</td>
852 <td class="paramname"><em>src_step_z</em>, </td>
853 </tr>
854 <tr>
855 <td class="paramkey"></td>
856 <td></td>
857 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]858 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
859 </tr>
860 <tr>
861 <td class="paramkey"></td>
862 <td></td>
863 <td class="paramtype">__global uchar *&#160;</td>
864 <td class="paramname"><em>max_ptr</em>, </td>
865 </tr>
866 <tr>
867 <td class="paramkey"></td>
868 <td></td>
869 <td class="paramtype">uint&#160;</td>
870 <td class="paramname"><em>max_stride_x</em>, </td>
871 </tr>
872 <tr>
873 <td class="paramkey"></td>
874 <td></td>
875 <td class="paramtype">uint&#160;</td>
876 <td class="paramname"><em>max_step_x</em>, </td>
877 </tr>
878 <tr>
879 <td class="paramkey"></td>
880 <td></td>
881 <td class="paramtype">uint&#160;</td>
882 <td class="paramname"><em>max_stride_y</em>, </td>
883 </tr>
884 <tr>
885 <td class="paramkey"></td>
886 <td></td>
887 <td class="paramtype">uint&#160;</td>
888 <td class="paramname"><em>max_step_y</em>, </td>
889 </tr>
890 <tr>
891 <td class="paramkey"></td>
892 <td></td>
893 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]894 <td class="paramname"><em>max_stride_z</em>, </td>
895 </tr>
896 <tr>
897 <td class="paramkey"></td>
898 <td></td>
899 <td class="paramtype">uint&#160;</td>
900 <td class="paramname"><em>max_step_z</em>, </td>
901 </tr>
902 <tr>
903 <td class="paramkey"></td>
904 <td></td>
905 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]906 <td class="paramname"><em>max_offset_first_element_in_bytes</em>, </td>
907 </tr>
908 <tr>
909 <td class="paramkey"></td>
910 <td></td>
911 <td class="paramtype">__global uchar *&#160;</td>
912 <td class="paramname"><em>dst_ptr</em>, </td>
913 </tr>
914 <tr>
915 <td class="paramkey"></td>
916 <td></td>
917 <td class="paramtype">uint&#160;</td>
918 <td class="paramname"><em>dst_stride_x</em>, </td>
919 </tr>
920 <tr>
921 <td class="paramkey"></td>
922 <td></td>
923 <td class="paramtype">uint&#160;</td>
924 <td class="paramname"><em>dst_step_x</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>dst_stride_y</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>dst_step_y</em>, </td>
937 </tr>
938 <tr>
939 <td class="paramkey"></td>
940 <td></td>
941 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]942 <td class="paramname"><em>dst_stride_z</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>dst_step_z</em>, </td>
949 </tr>
950 <tr>
951 <td class="paramkey"></td>
952 <td></td>
953 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]954 <td class="paramname"><em>dst_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>sum_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>sum_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>sum_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>sum_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>sum_step_y</em>, </td>
985 </tr>
986 <tr>
987 <td class="paramkey"></td>
988 <td></td>
989 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]990 <td class="paramname"><em>sum_stride_z</em>, </td>
991 </tr>
992 <tr>
993 <td class="paramkey"></td>
994 <td></td>
995 <td class="paramtype">uint&#160;</td>
996 <td class="paramname"><em>sum_step_z</em>, </td>
997 </tr>
998 <tr>
999 <td class="paramkey"></td>
1000 <td></td>
1001 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1002 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1003 </tr>
1004 <tr>
1005 <td class="paramkey"></td>
1006 <td></td>
1007 <td class="paramtype">uint&#160;</td>
1008 <td class="paramname"><em>width</em>&#160;</td>
1009 </tr>
1010 <tr>
1011 <td></td>
1012 <td>)</td>
1013 <td></td><td></td>
1014 </tr>
1015 </table>
1016</div><div class="memdoc">
1017
1018<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>
1019<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>
1020<dd>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1021Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1022<dd>
1023In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed.</dd></dl>
1024<dl class="params"><dt>Parameters</dt><dd>
1025 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1026 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: QS8/QS16/F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1027 <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>
1028 <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>
1029 <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>
1030 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1031 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
1032 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1033 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1034 <tr><td class="paramdir">[in]</td><td class="paramname">max_ptr</td><td>Pointer to the max values tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1035 <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>
1036 <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>
1037 <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>
1038 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1039 <tr><td class="paramdir">[in]</td><td class="paramname">max_stride_z</td><td>Stride of the max values tensor in Z dimension (in bytes) </td></tr>
1040 <tr><td class="paramdir">[in]</td><td class="paramname">max_step_z</td><td>max_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1041 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1042 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1043 <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>
1044 <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>
1045 <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>
1046 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1047 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
1048 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1049 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1050 <tr><td class="paramdir">[out]</td><td class="paramname">sum_ptr</td><td>Pointer to the sum values tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1051 <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>
1052 <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>
1053 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1054 <tr><td class="paramdir">[in]</td><td class="paramname">sum_step_y</td><td>sum_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
1055 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_z</td><td>Stride of the sum values tensor in Z dimension (in bytes) </td></tr>
1056 <tr><td class="paramdir">[in]</td><td class="paramname">sum_step_z</td><td>sum_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1057 <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>
1058 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1059 </table>
1060 </dd>
1061</dl>
1062
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1063<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00166">166</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1064
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1065<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">ADD_OP</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00105">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00046">EXP_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00059">idx16</a>, <a class="el" href="tests_2validation_2_fixed_point_8h_source.xhtml#l00889">arm_compute::test::fixed_point_arithmetic::detail::max()</a>, <a class="el" href="helpers_8h_source.xhtml#l00292">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00133">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00053">SELECT_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, and <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>.</p>
1066<div class="fragment"><div class="line"><a name="l00172"></a><span class="lineno"> 172</span>&#160;{</div>
1067<div class="line"><a name="l00173"></a><span class="lineno"> 173</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div>
1068<div class="line"><a name="l00174"></a><span class="lineno"> 174</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> dst = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div>
1069<div class="line"><a name="l00175"></a><span class="lineno"> 175</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">max</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(max);</div>
1070<div class="line"><a name="l00176"></a><span class="lineno"> 176</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(sum);</div>
1071<div class="line"><a name="l00177"></a><span class="lineno"> 177</span>&#160;</div>
1072<div class="line"><a name="l00178"></a><span class="lineno"> 178</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div>
1073<div class="line"><a name="l00179"></a><span class="lineno"> 179</span>&#160; <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> max_val = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;max, 0, 0));</div>
1074<div class="line"><a name="l00180"></a><span class="lineno"> 180</span>&#160;</div>
1075<div class="line"><a name="l00181"></a><span class="lineno"> 181</span>&#160; <span class="comment">// Set sum vector</span></div>
1076<div class="line"><a name="l00182"></a><span class="lineno"> 182</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div>
1077<div class="line"><a name="l00183"></a><span class="lineno"> 183</span>&#160; sum1D = 0;</div>
1078<div class="line"><a name="l00184"></a><span class="lineno"> 184</span>&#160;</div>
1079<div class="line"><a name="l00185"></a><span class="lineno"> 185</span>&#160; <span class="comment">// Shift values, exp and sum</span></div>
1080<div class="line"><a name="l00186"></a><span class="lineno"> 186</span>&#160; const uint width4 = width &gt;&gt; 4;</div>
1081<div class="line"><a name="l00187"></a><span class="lineno"> 187</span>&#160; for(uint i = 0; i &lt; width4; i++)</div>
1082<div class="line"><a name="l00188"></a><span class="lineno"> 188</span>&#160; {</div>
1083<div class="line"><a name="l00189"></a><span class="lineno"> 189</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div>
1084<div class="line"><a name="l00190"></a><span class="lineno"> 190</span>&#160; data = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; 4, 0));</div>
1085<div class="line"><a name="l00191"></a><span class="lineno"> 191</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16);</div>
1086<div class="line"><a name="l00192"></a><span class="lineno"> 192</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, 16);</div>
1087<div class="line"><a name="l00193"></a><span class="lineno"> 193</span>&#160; vstore16(data, 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, i &lt;&lt; 4, 0));</div>
1088<div class="line"><a name="l00194"></a><span class="lineno"> 194</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, 16);</div>
1089<div class="line"><a name="l00195"></a><span class="lineno"> 195</span>&#160; }</div>
1090<div class="line"><a name="l00196"></a><span class="lineno"> 196</span>&#160;</div>
1091<div class="line"><a name="l00197"></a><span class="lineno"> 197</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div>
1092<div class="line"><a name="l00198"></a><span class="lineno"> 198</span>&#160;<span class="preprocessor"></span> <span class="comment">// Handle non multiple of 16</span></div>
1093<div class="line"><a name="l00199"></a><span class="lineno"> 199</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div>
1094<div class="line"><a name="l00200"></a><span class="lineno"> 200</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>
1095<div class="line"><a name="l00201"></a><span class="lineno"> 201</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, 16);</div>
1096<div class="line"><a name="l00202"></a><span class="lineno"> 202</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, 16);</div>
1097<div class="line"><a name="l00203"></a><span class="lineno"> 203</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 16)</div>
1098<div class="line"><a name="l00204"></a><span class="lineno"> 204</span>&#160; widx = <a class="code" href="fixed__point_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="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 16));</div>
1099<div class="line"><a name="l00205"></a><span class="lineno"> 205</span>&#160; data = select(0, data, widx);</div>
1100<div class="line"><a name="l00206"></a><span class="lineno"> 206</span>&#160; vstore16(data, 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, width4 &lt;&lt; 4, 0));</div>
1101<div class="line"><a name="l00207"></a><span class="lineno"> 207</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, 16);</div>
1102<div class="line"><a name="l00208"></a><span class="lineno"> 208</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_16 */</span><span class="preprocessor"></span></div>
1103<div class="line"><a name="l00209"></a><span class="lineno"> 209</span>&#160;<span class="preprocessor"></span></div>
1104<div class="line"><a name="l00210"></a><span class="lineno"> 210</span>&#160; <span class="comment">// Perform min/max reduction</span></div>
1105<div class="line"><a name="l00211"></a><span class="lineno"> 211</span>&#160; sum1D.s01234567 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01234567, sum1D.s89ABCDEF, DATA_TYPE, 8);</div>
1106<div class="line"><a name="l00212"></a><span class="lineno"> 212</span>&#160; sum1D.s0123 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0123, sum1D.s4567, DATA_TYPE, 4);</div>
1107<div class="line"><a name="l00213"></a><span class="lineno"> 213</span>&#160; sum1D.s01 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01, sum1D.s23, DATA_TYPE, 2);</div>
1108<div class="line"><a name="l00214"></a><span class="lineno"> 214</span>&#160; sum1D.s0 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0, sum1D.s1, DATA_TYPE, 1);</div>
1109<div class="line"><a name="l00215"></a><span class="lineno"> 215</span>&#160;</div>
1110<div class="line"><a name="l00216"></a><span class="lineno"> 216</span>&#160; <span class="comment">// Calculate and store result</span></div>
1111<div class="line"><a name="l00217"></a><span class="lineno"> 217</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div>
1112<div class="line"><a name="l00218"></a><span class="lineno"> 218</span>&#160;}</div>
1113<div class="ttc" id="fixed__point_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00098">fixed_point.h:98</a></div></div>
1114<div class="ttc" id="softmax__layer_8cl_xhtml_a93cf800667317d96574477b9f0a75234"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a></div><div class="ttdeci">#define EXP_OP(x, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00046">softmax_layer.cl:46</a></div></div>
1115<div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
1116<div class="ttc" id="reduction__operation_8cl_xhtml_ab0df00f5333da51860deb93deb44a782"><div class="ttname"><a href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a></div><div class="ttdeci">DATA_TYPE sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate sum of a vector. </div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00052">reduction_operation.cl:52</a></div></div>
1117<div class="ttc" id="softmax__layer_8cl_xhtml_a44206a4e5783c7aabacec88aad878c88"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a></div><div class="ttdeci">#define ADD_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00043">softmax_layer.cl:43</a></div></div>
1118<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#l00292">helpers.h:292</a></div></div>
1119<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#l00053">softmax_layer.cl:53</a></div></div>
1120<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#l00131">helpers.h:131</a></div></div>
1121<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#l00133">helpers.h:133</a></div></div>
1122<div class="ttc" id="softmax__layer_8cl_xhtml_ac3af2d18008cbbf7247ae48fcd6e0c4e"><div class="ttname"><a href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a></div><div class="ttdeci">#define SUB_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00044">softmax_layer.cl:44</a></div></div>
1123<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_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="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
1124<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#l00059">softmax_layer.cl:59</a></div></div>
1125<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00105">helpers.h:105</a></div></div>
1126<div class="ttc" id="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail_xhtml_ad91bb73431b4de1f4946ed949d444849"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">arm_compute::test::fixed_point_arithmetic::detail::max</a></div><div class="ttdeci">fixed_point&lt; T &gt; max(fixed_point&lt; T &gt; x, fixed_point&lt; T &gt; y)</div><div class="ttdef"><b>Definition:</b> <a href="tests_2validation_2_fixed_point_8h_source.xhtml#l00889">FixedPoint.h:889</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1127</div><!-- fragment -->
1128</div>
1129</div>
1130<h2 class="groupheader">Variable Documentation</h2>
1131<a class="anchor" id="a0712735973f172ac9efc7d48a31e47ad"></a>
1132<div class="memitem">
1133<div class="memproto">
1134 <table class="memname">
1135 <tr>
1136 <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1137 </tr>
1138 </table>
1139</div><div class="memdoc">
1140
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1141<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1142
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1143<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00085">softmax_layer_max()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00166">softmax_layer_shift_exp_sum()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1144
1145</div>
1146</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1147<a class="anchor" id="a538b4b63f40e7b12891774e03a4f0dec"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1148<div class="memitem">
1149<div class="memproto">
1150 <table class="memname">
1151 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1152 <td class="memname">__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( -FLT_MAX )</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1153 </tr>
1154 </table>
1155</div><div class="memdoc">
1156
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1157<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00058">58</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1158
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1159<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00085">softmax_layer_max()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1160
1161</div>
1162</div>
1163</div><!-- contents -->
1164</div><!-- doc-content -->
1165<!-- start footer part -->
1166<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
1167 <ul>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1168 <li class="navelem"><a class="el" href="dir_55b5a5006e943fb664ff8cff3cfe7768.xhtml">src</a></li><li class="navelem"><a class="el" href="dir_f6c3ae42d7e27145d0a3c3e38ca1c000.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_1b7ae9123c88e650426d50d95c19d414.xhtml">CL</a></li><li class="navelem"><a class="el" href="dir_b7c3d80b0dadb2139bc73b7161751c12.xhtml">cl_kernels</a></li><li class="navelem"><a class="el" href="softmax__layer_8cl.xhtml">softmax_layer.cl</a></li>
1169 <li class="footer">Generated on Thu Sep 28 2017 14:37:54 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1170 <a href="http://www.doxygen.org/index.html">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame^]1171 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.6 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1172 </ul>
1173</div>
1174</body>
1175</html>