Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 292227986edb37b01061afcad6df18ba9d6ccbeb / . / documentation / softmax__layer_8cl.xhtml
blob: e9579f345bbf0dc9110d8e6215d13196b0879f2d [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"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]7<meta name="generator" content="Doxygen 1.8.11"/>
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 Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]15<script type="text/javascript" src="navtreedata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]16<script type="text/javascript" src="navtree.js"></script>
17<script type="text/javascript">
18 $(document).ready(initResizable);
19 $(window).load(resizeHeight);
20</script>
21<link href="search/search.css" rel="stylesheet" type="text/css"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]22<script type="text/javascript" src="search/searchdata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]23<script type="text/javascript" src="search/search.js"></script>
24<script type="text/javascript">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]25 $(document).ready(function() { init_search(); });
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]26</script>
27<script type="text/x-mathjax-config">
28 MathJax.Hub.Config({
29 extensions: ["tex2jax.js"],
30 jax: ["input/TeX","output/HTML-CSS"],
31});
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]32</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]33<link href="doxygen.css" rel="stylesheet" type="text/css" />
34</head>
35<body>
36<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
37<div id="titlearea">
38<table cellspacing="0" cellpadding="0">
39 <tbody>
40 <tr style="height: 56px;">
41 <td style="padding-left: 0.5em;">
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]42 <div id="projectname">Compute Library
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]43 &#160;<span id="projectnumber">18.01</span>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]44 </div>
45 </td>
46 </tr>
47 </tbody>
48</table>
49</div>
50<!-- end header part -->
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]51<!-- Generated by Doxygen 1.8.11 -->
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]52<script type="text/javascript">
53var searchBox = new SearchBox("searchBox", "search",false,'Search');
54</script>
55 <div id="navrow1" class="tabs">
56 <ul class="tablist">
57 <li><a href="index.xhtml"><span>Main&#160;Page</span></a></li>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]58 <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]59 <li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
60 <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
61 <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
62 <li>
63 <div id="MSearchBox" class="MSearchBoxInactive">
64 <span class="left">
65 <img id="MSearchSelect" src="search/mag_sel.png"
66 onmouseover="return searchBox.OnSearchSelectShow()"
67 onmouseout="return searchBox.OnSearchSelectHide()"
68 alt=""/>
69 <input type="text" id="MSearchField" value="Search" accesskey="S"
70 onfocus="searchBox.OnSearchFieldFocus(true)"
71 onblur="searchBox.OnSearchFieldFocus(false)"
72 onkeyup="searchBox.OnSearchFieldChange(event)"/>
73 </span><span class="right">
74 <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
75 </span>
76 </div>
77 </li>
78 </ul>
79 </div>
80 <div id="navrow2" class="tabs2">
81 <ul class="tablist">
82 <li><a href="files.xhtml"><span>File&#160;List</span></a></li>
83 <li><a href="globals.xhtml"><span>Globals</span></a></li>
84 </ul>
85 </div>
86</div><!-- top -->
87<div id="side-nav" class="ui-resizable side-nav-resizable">
88 <div id="nav-tree">
89 <div id="nav-tree-contents">
90 <div id="nav-sync" class="sync"></div>
91 </div>
92 </div>
93 <div id="splitbar" style="-moz-user-select:none;"
94 class="ui-resizable-handle">
95 </div>
96</div>
97<script type="text/javascript">
98$(document).ready(function(){initNavTree('softmax__layer_8cl.xhtml','');});
99</script>
100<div id="doc-content">
101<!-- window showing the filter options -->
102<div id="MSearchSelectWindow"
103 onmouseover="return searchBox.OnSearchSelectShow()"
104 onmouseout="return searchBox.OnSearchSelectHide()"
105 onkeydown="return searchBox.OnSearchSelectKey(event)">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]106</div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]107
108<!-- iframe showing the search results (closed by default) -->
109<div id="MSearchResultsWindow">
110<iframe src="javascript:void(0)" frameborder="0"
111 name="MSearchResults" id="MSearchResults">
112</iframe>
113</div>
114
115<div class="header">
116 <div class="summary">
117<a href="#define-members">Macros</a> &#124;
118<a href="#func-members">Functions</a> &#124;
119<a href="#var-members">Variables</a> </div>
120 <div class="headertitle">
121<div class="title">softmax_layer.cl File Reference</div> </div>
122</div><!--header-->
123<div class="contents">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]124<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]125</div>
126<p><a href="softmax__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
127<table class="memberdecls">
128<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
129Macros</h2></td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]130<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]131<tr class="separator:abaa48ad818c44e415fd3f9dd0f27bf01"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]132<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]133<tr class="separator:a44206a4e5783c7aabacec88aad878c88"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]134<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]135<tr class="separator:ac3af2d18008cbbf7247ae48fcd6e0c4e"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]136<tr class="memitem:a22303c4047ec5027c1538d53964b9d0d"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(x, y, type, <a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>)&#160;&#160;&#160;((x) * (y))</td></tr>
137<tr class="separator:a22303c4047ec5027c1538d53964b9d0d"><td class="memSeparator" colspan="2">&#160;</td></tr>
138<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]139<tr class="separator:a8cde99b1ce0f3c1dacd49261b0cf03d8"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]140<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]141<tr class="separator:a93cf800667317d96574477b9f0a75234"><td class="memSeparator" colspan="2">&#160;</td></tr>
142<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]143<tr class="separator:a80b22c555ddadb47cc6ca338a9c49126"><td class="memSeparator" colspan="2">&#160;</td></tr>
144<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>
145<tr class="separator:af5b2e33e3c5fcaab3a213f26c2300170"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]146<tr class="memitem:a08246606c233e7785a497c09672f366f"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>&#160;&#160;&#160;1</td></tr>
147<tr class="separator:a08246606c233e7785a497c09672f366f"><td class="memSeparator" colspan="2">&#160;</td></tr>
148<tr class="memitem:a7c78836761fa3b5b124efea237dac70f"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>&#160;&#160;&#160;16</td></tr>
149<tr class="separator:a7c78836761fa3b5b124efea237dac70f"><td class="memSeparator" colspan="2">&#160;</td></tr>
150<tr class="memitem:a372393c380805985b813dbb16d589a64"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>&#160;&#160;&#160;4</td></tr>
151<tr class="separator:a372393c380805985b813dbb16d589a64"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]152</table><table class="memberdecls">
153<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
154Functions</h2></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]155<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]156<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]157<tr class="separator:a9e86f72ca0d249cb88c36bd50620da8a"><td class="memSeparator" colspan="2">&#160;</td></tr>
158<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]159<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]160<tr class="separator:aa4412bb319c67deaeab07a481d93dfdb"><td class="memSeparator" colspan="2">&#160;</td></tr>
161<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]162<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]163<tr class="separator:ac4247ac0991e85965b7ded764e78f12c"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]164<tr class="memitem:a9d74601bcabbb2f14bcf52385ad666dc"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a9d74601bcabbb2f14bcf52385ad666dc">softmax_layer_max_shift_exp_sum_serial</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 *maxo_ptr, uint maxo_stride_x, uint maxo_step_x, uint maxo_stride_y, uint maxo_step_y, uint maxo_stride_z, uint maxo_step_z, uint maxo_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>
165<tr class="memdesc:a9d74601bcabbb2f14bcf52385ad666dc"><td class="mdescLeft">&#160;</td><td class="mdescRight">Identifies the maximum value across the 1st dimension and shifts the values of the input tensor by this maximum value, then gets the exponent of each element as sums all elements across each row. <a href="#a9d74601bcabbb2f14bcf52385ad666dc">More...</a><br /></td></tr>
166<tr class="separator:a9d74601bcabbb2f14bcf52385ad666dc"><td class="memSeparator" colspan="2">&#160;</td></tr>
167<tr class="memitem:a82bc360279fcf7cf8033dca6022f21f5"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a82bc360279fcf7cf8033dca6022f21f5">softmax_layer_max_shift_exp_sum_parallel</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 *maxo_ptr, uint maxo_stride_x, uint maxo_step_x, uint maxo_stride_y, uint maxo_step_y, uint maxo_stride_z, uint maxo_step_z, uint maxo_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>
168<tr class="memdesc:a82bc360279fcf7cf8033dca6022f21f5"><td class="mdescLeft">&#160;</td><td class="mdescRight">Identifies the maximum value across the 1st dimension and shifts the values of the input tensor by this maximum value, then gets the exponent of each element as sums all elements across each row. <a href="#a82bc360279fcf7cf8033dca6022f21f5">More...</a><br /></td></tr>
169<tr class="separator:a82bc360279fcf7cf8033dca6022f21f5"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]170</table><table class="memberdecls">
171<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="var-members"></a>
172Variables</h2></td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]173<tr class="memitem:af7a08044d0e491a0ee1520a24a107a2b"><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#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a> = ( DATA_TYPE16 )( -FLT_MAX )</td></tr>
174<tr class="separator:af7a08044d0e491a0ee1520a24a107a2b"><td class="memSeparator" colspan="2">&#160;</td></tr>
175<tr class="memitem:aa1dd94b8d98f1c6d790bdf0fc5de29e9"><td class="memItemLeft" align="right" valign="top">__constant uint16&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a> = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td></tr>
176<tr class="separator:aa1dd94b8d98f1c6d790bdf0fc5de29e9"><td class="memSeparator" colspan="2">&#160;</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]177<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>
178<tr class="separator:a538b4b63f40e7b12891774e03a4f0dec"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]179<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>
180<tr class="separator:a0712735973f172ac9efc7d48a31e47ad"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]181<tr class="memitem:a4884a666a1e93fbf8c27bd7d2da3c8bb"><td class="memItemLeft" align="right" valign="top">__constant uint4&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a> = (uint4)(0, 1, 2, 3)</td></tr>
182<tr class="separator:a4884a666a1e93fbf8c27bd7d2da3c8bb"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]183</table>
184<h2 class="groupheader">Macro Definition Documentation</h2>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]185<a class="anchor" id="a44206a4e5783c7aabacec88aad878c88"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]186<div class="memitem">
187<div class="memproto">
188 <table class="memname">
189 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]190 <td class="memname">#define ADD_OP</td>
191 <td>(</td>
192 <td class="paramtype">&#160;</td>
193 <td class="paramname">x, </td>
194 </tr>
195 <tr>
196 <td class="paramkey"></td>
197 <td></td>
198 <td class="paramtype">&#160;</td>
199 <td class="paramname">y, </td>
200 </tr>
201 <tr>
202 <td class="paramkey"></td>
203 <td></td>
204 <td class="paramtype">&#160;</td>
205 <td class="paramname">type, </td>
206 </tr>
207 <tr>
208 <td class="paramkey"></td>
209 <td></td>
210 <td class="paramtype">&#160;</td>
211 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
212 </tr>
213 <tr>
214 <td></td>
215 <td>)</td>
216 <td></td><td>&#160;&#160;&#160;((x) + (y))</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]217 </tr>
218 </table>
219</div><div class="memdoc">
220
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]221<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]222
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]223<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00196">softmax_layer_shift_exp_sum()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]224
225</div>
226</div>
227<a class="anchor" id="a8cde99b1ce0f3c1dacd49261b0cf03d8"></a>
228<div class="memitem">
229<div class="memproto">
230 <table class="memname">
231 <tr>
232 <td class="memname">#define DIV_OP</td>
233 <td>(</td>
234 <td class="paramtype">&#160;</td>
235 <td class="paramname">x, </td>
236 </tr>
237 <tr>
238 <td class="paramkey"></td>
239 <td></td>
240 <td class="paramtype">&#160;</td>
241 <td class="paramname">y, </td>
242 </tr>
243 <tr>
244 <td class="paramkey"></td>
245 <td></td>
246 <td class="paramtype">&#160;</td>
247 <td class="paramname">type, </td>
248 </tr>
249 <tr>
250 <td class="paramkey"></td>
251 <td></td>
252 <td class="paramtype">&#160;</td>
253 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
254 </tr>
255 <tr>
256 <td></td>
257 <td>)</td>
258 <td></td><td>&#160;&#160;&#160;((x) / (y))</td>
259 </tr>
260 </table>
261</div><div class="memdoc">
262
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]263<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00047">47</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]264
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]265<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00292">softmax_layer_norm()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]266
267</div>
268</div>
269<a class="anchor" id="a93cf800667317d96574477b9f0a75234"></a>
270<div class="memitem">
271<div class="memproto">
272 <table class="memname">
273 <tr>
274 <td class="memname">#define EXP_OP</td>
275 <td>(</td>
276 <td class="paramtype">&#160;</td>
277 <td class="paramname">x, </td>
278 </tr>
279 <tr>
280 <td class="paramkey"></td>
281 <td></td>
282 <td class="paramtype">&#160;</td>
283 <td class="paramname">type, </td>
284 </tr>
285 <tr>
286 <td class="paramkey"></td>
287 <td></td>
288 <td class="paramtype">&#160;</td>
289 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
290 </tr>
291 <tr>
292 <td></td>
293 <td>)</td>
294 <td></td><td>&#160;&#160;&#160;exp((x))</td>
295 </tr>
296 </table>
297</div><div class="memdoc">
298
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]299<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00048">48</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]300
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]301<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00196">softmax_layer_shift_exp_sum()</a>.</p>
302
303</div>
304</div>
305<a class="anchor" id="a08246606c233e7785a497c09672f366f"></a>
306<div class="memitem">
307<div class="memproto">
308 <table class="memname">
309 <tr>
310 <td class="memname">#define GRID_SIZE&#160;&#160;&#160;1</td>
311 </tr>
312 </table>
313</div><div class="memdoc">
314
315<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">62</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
316
317<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>.</p>
318
319</div>
320</div>
321<a class="anchor" id="a372393c380805985b813dbb16d589a64"></a>
322<div class="memitem">
323<div class="memproto">
324 <table class="memname">
325 <tr>
326 <td class="memname">#define LOG_VECTOR_SIZE&#160;&#160;&#160;4</td>
327 </tr>
328 </table>
329</div><div class="memdoc">
330
331<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00080">80</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
332
333<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]334
335</div>
336</div>
337<a class="anchor" id="abaa48ad818c44e415fd3f9dd0f27bf01"></a>
338<div class="memitem">
339<div class="memproto">
340 <table class="memname">
341 <tr>
342 <td class="memname">#define MAX_OP</td>
343 <td>(</td>
344 <td class="paramtype">&#160;</td>
345 <td class="paramname">x, </td>
346 </tr>
347 <tr>
348 <td class="paramkey"></td>
349 <td></td>
350 <td class="paramtype">&#160;</td>
351 <td class="paramname">y, </td>
352 </tr>
353 <tr>
354 <td class="paramkey"></td>
355 <td></td>
356 <td class="paramtype">&#160;</td>
357 <td class="paramname">type, </td>
358 </tr>
359 <tr>
360 <td class="paramkey"></td>
361 <td></td>
362 <td class="paramtype">&#160;</td>
363 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
364 </tr>
365 <tr>
366 <td></td>
367 <td>)</td>
368 <td></td><td>&#160;&#160;&#160;max((x), (y))</td>
369 </tr>
370 </table>
371</div><div class="memdoc">
372
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]373<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]374
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]375<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00114">softmax_layer_max()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]376
377</div>
378</div>
379<a class="anchor" id="a80b22c555ddadb47cc6ca338a9c49126"></a>
380<div class="memitem">
381<div class="memproto">
382 <table class="memname">
383 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]384 <td class="memname">#define MINVAL&#160;&#160;&#160;-FLT_MAX</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]385 </tr>
386 </table>
387</div><div class="memdoc">
388
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]389<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00054">54</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
390
391</div>
392</div>
393<a class="anchor" id="a22303c4047ec5027c1538d53964b9d0d"></a>
394<div class="memitem">
395<div class="memproto">
396 <table class="memname">
397 <tr>
398 <td class="memname">#define MUL_OP</td>
399 <td>(</td>
400 <td class="paramtype">&#160;</td>
401 <td class="paramname">x, </td>
402 </tr>
403 <tr>
404 <td class="paramkey"></td>
405 <td></td>
406 <td class="paramtype">&#160;</td>
407 <td class="paramname">y, </td>
408 </tr>
409 <tr>
410 <td class="paramkey"></td>
411 <td></td>
412 <td class="paramtype">&#160;</td>
413 <td class="paramname">type, </td>
414 </tr>
415 <tr>
416 <td class="paramkey"></td>
417 <td></td>
418 <td class="paramtype">&#160;</td>
419 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
420 </tr>
421 <tr>
422 <td></td>
423 <td>)</td>
424 <td></td><td>&#160;&#160;&#160;((x) * (y))</td>
425 </tr>
426 </table>
427</div><div class="memdoc">
428
429<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>
430
431<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00196">softmax_layer_shift_exp_sum()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]432
433</div>
434</div>
435<a class="anchor" id="af5b2e33e3c5fcaab3a213f26c2300170"></a>
436<div class="memitem">
437<div class="memproto">
438 <table class="memname">
439 <tr>
440 <td class="memname">#define SELECT_DATA_TYPE&#160;&#160;&#160;int</td>
441 </tr>
442 </table>
443</div><div class="memdoc">
444
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]445<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">55</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]446
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]447<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00114">softmax_layer_max()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00196">softmax_layer_shift_exp_sum()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]448
449</div>
450</div>
451<a class="anchor" id="ac3af2d18008cbbf7247ae48fcd6e0c4e"></a>
452<div class="memitem">
453<div class="memproto">
454 <table class="memname">
455 <tr>
456 <td class="memname">#define SUB_OP</td>
457 <td>(</td>
458 <td class="paramtype">&#160;</td>
459 <td class="paramname">x, </td>
460 </tr>
461 <tr>
462 <td class="paramkey"></td>
463 <td></td>
464 <td class="paramtype">&#160;</td>
465 <td class="paramname">y, </td>
466 </tr>
467 <tr>
468 <td class="paramkey"></td>
469 <td></td>
470 <td class="paramtype">&#160;</td>
471 <td class="paramname">type, </td>
472 </tr>
473 <tr>
474 <td class="paramkey"></td>
475 <td></td>
476 <td class="paramtype">&#160;</td>
477 <td class="paramname"><a class="el" href="hwc_8hpp.xhtml#ab2c6b258f02add8fdf4cfc7c371dd772">size</a>&#160;</td>
478 </tr>
479 <tr>
480 <td></td>
481 <td>)</td>
482 <td></td><td>&#160;&#160;&#160;((x) - (y))</td>
483 </tr>
484 </table>
485</div><div class="memdoc">
486
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]487<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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]488
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]489<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00196">softmax_layer_shift_exp_sum()</a>.</p>
490
491</div>
492</div>
493<a class="anchor" id="a7c78836761fa3b5b124efea237dac70f"></a>
494<div class="memitem">
495<div class="memproto">
496 <table class="memname">
497 <tr>
498 <td class="memname">#define VECTOR_SIZE&#160;&#160;&#160;16</td>
499 </tr>
500 </table>
501</div><div class="memdoc">
502
503<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00079">79</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
504
505<p>Referenced by <a class="el" href="gemm_8cl_source.xhtml#l00223">gemm_interleave4x4_8bit()</a>, <a class="el" href="convolution__layer_8cl_source.xhtml#l00306">im2col_reduced()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]506
507</div>
508</div>
509<h2 class="groupheader">Function Documentation</h2>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]510<a class="anchor" id="a9e86f72ca0d249cb88c36bd50620da8a"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]511<div class="memitem">
512<div class="memproto">
513 <table class="memname">
514 <tr>
515 <td class="memname">__kernel void softmax_layer_max </td>
516 <td>(</td>
517 <td class="paramtype">__global uchar *&#160;</td>
518 <td class="paramname"><em>src_ptr</em>, </td>
519 </tr>
520 <tr>
521 <td class="paramkey"></td>
522 <td></td>
523 <td class="paramtype">uint&#160;</td>
524 <td class="paramname"><em>src_stride_x</em>, </td>
525 </tr>
526 <tr>
527 <td class="paramkey"></td>
528 <td></td>
529 <td class="paramtype">uint&#160;</td>
530 <td class="paramname"><em>src_step_x</em>, </td>
531 </tr>
532 <tr>
533 <td class="paramkey"></td>
534 <td></td>
535 <td class="paramtype">uint&#160;</td>
536 <td class="paramname"><em>src_stride_y</em>, </td>
537 </tr>
538 <tr>
539 <td class="paramkey"></td>
540 <td></td>
541 <td class="paramtype">uint&#160;</td>
542 <td class="paramname"><em>src_step_y</em>, </td>
543 </tr>
544 <tr>
545 <td class="paramkey"></td>
546 <td></td>
547 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]548 <td class="paramname"><em>src_stride_z</em>, </td>
549 </tr>
550 <tr>
551 <td class="paramkey"></td>
552 <td></td>
553 <td class="paramtype">uint&#160;</td>
554 <td class="paramname"><em>src_step_z</em>, </td>
555 </tr>
556 <tr>
557 <td class="paramkey"></td>
558 <td></td>
559 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]560 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
561 </tr>
562 <tr>
563 <td class="paramkey"></td>
564 <td></td>
565 <td class="paramtype">__global uchar *&#160;</td>
566 <td class="paramname"><em>dst_ptr</em>, </td>
567 </tr>
568 <tr>
569 <td class="paramkey"></td>
570 <td></td>
571 <td class="paramtype">uint&#160;</td>
572 <td class="paramname"><em>dst_stride_x</em>, </td>
573 </tr>
574 <tr>
575 <td class="paramkey"></td>
576 <td></td>
577 <td class="paramtype">uint&#160;</td>
578 <td class="paramname"><em>dst_step_x</em>, </td>
579 </tr>
580 <tr>
581 <td class="paramkey"></td>
582 <td></td>
583 <td class="paramtype">uint&#160;</td>
584 <td class="paramname"><em>dst_stride_y</em>, </td>
585 </tr>
586 <tr>
587 <td class="paramkey"></td>
588 <td></td>
589 <td class="paramtype">uint&#160;</td>
590 <td class="paramname"><em>dst_step_y</em>, </td>
591 </tr>
592 <tr>
593 <td class="paramkey"></td>
594 <td></td>
595 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]596 <td class="paramname"><em>dst_stride_z</em>, </td>
597 </tr>
598 <tr>
599 <td class="paramkey"></td>
600 <td></td>
601 <td class="paramtype">uint&#160;</td>
602 <td class="paramname"><em>dst_step_z</em>, </td>
603 </tr>
604 <tr>
605 <td class="paramkey"></td>
606 <td></td>
607 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]608 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
609 </tr>
610 <tr>
611 <td class="paramkey"></td>
612 <td></td>
613 <td class="paramtype">uint&#160;</td>
614 <td class="paramname"><em>width</em>&#160;</td>
615 </tr>
616 <tr>
617 <td></td>
618 <td>)</td>
619 <td></td><td></td>
620 </tr>
621 </table>
622</div><div class="memdoc">
623
624<p>Identifies the maximum value across the 1st dimension. </p>
625<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>
626<dd>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]627Fixed 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]628<dd>
629In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed.</dd></dl>
630<dl class="params"><dt>Parameters</dt><dd>
631 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]632 <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]633 <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>
634 <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>
635 <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>
636 <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]637 <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>
638 <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]639 <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]640 <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]641 <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>
642 <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>
643 <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>
644 <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]645 <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>
646 <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]647 <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>
648 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
649 </table>
650 </dd>
651</dl>
652
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]653<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00114">114</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]654
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]655<p>References <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00087">idx16</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">MAX_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">SELECT_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00086">type_min</a>, and <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>.</p>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]656<div class="fragment"><div class="line"><a name="l00118"></a><span class="lineno"> 118</span>&#160;{</div><div class="line"><a name="l00119"></a><span class="lineno"> 119</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00120"></a><span class="lineno"> 120</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00121"></a><span class="lineno"> 121</span>&#160;</div><div class="line"><a name="l00122"></a><span class="lineno"> 122</span>&#160; <span class="comment">// Initialize local maximum</span></div><div class="line"><a name="l00123"></a><span class="lineno"> 123</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><div class="line"><a name="l00124"></a><span class="lineno"> 124</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><div class="line"><a name="l00125"></a><span class="lineno"> 125</span>&#160;</div><div class="line"><a name="l00126"></a><span class="lineno"> 126</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00127"></a><span class="lineno"> 127</span>&#160; const uint width4 = width &gt;&gt; 4;</div><div class="line"><a name="l00128"></a><span class="lineno"> 128</span>&#160; for(uint i = 0; i &lt; width4; i++)</div><div class="line"><a name="l00129"></a><span class="lineno"> 129</span>&#160; {</div><div class="line"><a name="l00130"></a><span class="lineno"> 130</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00131"></a><span class="lineno"> 131</span>&#160; data = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; 4, 0));</div><div class="line"><a name="l00132"></a><span class="lineno"> 132</span>&#160; max_val = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data, max_val, DATA_TYPE, 16);</div><div class="line"><a name="l00133"></a><span class="lineno"> 133</span>&#160; }</div><div class="line"><a name="l00134"></a><span class="lineno"> 134</span>&#160;</div><div class="line"><a name="l00135"></a><span class="lineno"> 135</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00136"></a><span class="lineno"> 136</span>&#160; <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00137"></a><span class="lineno"> 137</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00138"></a><span class="lineno"> 138</span>&#160; data = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width4 &lt;&lt; 4, 0));</div><div class="line"><a name="l00139"></a><span class="lineno"> 139</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><div class="line"><a name="l00140"></a><span class="lineno"> 140</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><div class="line"><a name="l00141"></a><span class="lineno"> 141</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><div class="line"><a name="l00142"></a><span class="lineno"> 142</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_16 */</span><span class="preprocessor"></span></div><div class="line"><a name="l00143"></a><span class="lineno"> 143</span>&#160;</div><div class="line"><a name="l00144"></a><span class="lineno"> 144</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00145"></a><span class="lineno"> 145</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><div class="line"><a name="l00146"></a><span class="lineno"> 146</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><div class="line"><a name="l00147"></a><span class="lineno"> 147</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><div class="line"><a name="l00148"></a><span class="lineno"> 148</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><div class="line"><a name="l00149"></a><span class="lineno"> 149</span>&#160;</div><div class="line"><a name="l00150"></a><span class="lineno"> 150</span>&#160; <span class="comment">// Store result</span></div><div class="line"><a name="l00151"></a><span class="lineno"> 151</span>&#160; *((__global DATA_TYPE *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = max_val.s0;</div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span>&#160;}</div><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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]657<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]658<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#l00043">softmax_layer.cl:43</a></div></div>
659<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
660<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#l00055">softmax_layer.cl:55</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]661<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00123">Convolution.cpp:123</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]662<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
663<div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00142">helpers.h:142</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]664<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]665<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#l00087">softmax_layer.cl:87</a></div></div>
666<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
667<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#l00086">softmax_layer.cl:86</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]668<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00133">Convolution.cpp:133</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]669</div><!-- fragment -->
670</div>
671</div>
672<a class="anchor" id="a82bc360279fcf7cf8033dca6022f21f5"></a>
673<div class="memitem">
674<div class="memproto">
675 <table class="memname">
676 <tr>
677 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_parallel </td>
678 <td>(</td>
679 <td class="paramtype">__global uchar *&#160;</td>
680 <td class="paramname"><em>src_ptr</em>, </td>
681 </tr>
682 <tr>
683 <td class="paramkey"></td>
684 <td></td>
685 <td class="paramtype">uint&#160;</td>
686 <td class="paramname"><em>src_stride_x</em>, </td>
687 </tr>
688 <tr>
689 <td class="paramkey"></td>
690 <td></td>
691 <td class="paramtype">uint&#160;</td>
692 <td class="paramname"><em>src_step_x</em>, </td>
693 </tr>
694 <tr>
695 <td class="paramkey"></td>
696 <td></td>
697 <td class="paramtype">uint&#160;</td>
698 <td class="paramname"><em>src_stride_y</em>, </td>
699 </tr>
700 <tr>
701 <td class="paramkey"></td>
702 <td></td>
703 <td class="paramtype">uint&#160;</td>
704 <td class="paramname"><em>src_step_y</em>, </td>
705 </tr>
706 <tr>
707 <td class="paramkey"></td>
708 <td></td>
709 <td class="paramtype">uint&#160;</td>
710 <td class="paramname"><em>src_stride_z</em>, </td>
711 </tr>
712 <tr>
713 <td class="paramkey"></td>
714 <td></td>
715 <td class="paramtype">uint&#160;</td>
716 <td class="paramname"><em>src_step_z</em>, </td>
717 </tr>
718 <tr>
719 <td class="paramkey"></td>
720 <td></td>
721 <td class="paramtype">uint&#160;</td>
722 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
723 </tr>
724 <tr>
725 <td class="paramkey"></td>
726 <td></td>
727 <td class="paramtype">__global uchar *&#160;</td>
728 <td class="paramname"><em>maxo_ptr</em>, </td>
729 </tr>
730 <tr>
731 <td class="paramkey"></td>
732 <td></td>
733 <td class="paramtype">uint&#160;</td>
734 <td class="paramname"><em>maxo_stride_x</em>, </td>
735 </tr>
736 <tr>
737 <td class="paramkey"></td>
738 <td></td>
739 <td class="paramtype">uint&#160;</td>
740 <td class="paramname"><em>maxo_step_x</em>, </td>
741 </tr>
742 <tr>
743 <td class="paramkey"></td>
744 <td></td>
745 <td class="paramtype">uint&#160;</td>
746 <td class="paramname"><em>maxo_stride_y</em>, </td>
747 </tr>
748 <tr>
749 <td class="paramkey"></td>
750 <td></td>
751 <td class="paramtype">uint&#160;</td>
752 <td class="paramname"><em>maxo_step_y</em>, </td>
753 </tr>
754 <tr>
755 <td class="paramkey"></td>
756 <td></td>
757 <td class="paramtype">uint&#160;</td>
758 <td class="paramname"><em>maxo_stride_z</em>, </td>
759 </tr>
760 <tr>
761 <td class="paramkey"></td>
762 <td></td>
763 <td class="paramtype">uint&#160;</td>
764 <td class="paramname"><em>maxo_step_z</em>, </td>
765 </tr>
766 <tr>
767 <td class="paramkey"></td>
768 <td></td>
769 <td class="paramtype">uint&#160;</td>
770 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
771 </tr>
772 <tr>
773 <td class="paramkey"></td>
774 <td></td>
775 <td class="paramtype">__global uchar *&#160;</td>
776 <td class="paramname"><em>dst_ptr</em>, </td>
777 </tr>
778 <tr>
779 <td class="paramkey"></td>
780 <td></td>
781 <td class="paramtype">uint&#160;</td>
782 <td class="paramname"><em>dst_stride_x</em>, </td>
783 </tr>
784 <tr>
785 <td class="paramkey"></td>
786 <td></td>
787 <td class="paramtype">uint&#160;</td>
788 <td class="paramname"><em>dst_step_x</em>, </td>
789 </tr>
790 <tr>
791 <td class="paramkey"></td>
792 <td></td>
793 <td class="paramtype">uint&#160;</td>
794 <td class="paramname"><em>dst_stride_y</em>, </td>
795 </tr>
796 <tr>
797 <td class="paramkey"></td>
798 <td></td>
799 <td class="paramtype">uint&#160;</td>
800 <td class="paramname"><em>dst_step_y</em>, </td>
801 </tr>
802 <tr>
803 <td class="paramkey"></td>
804 <td></td>
805 <td class="paramtype">uint&#160;</td>
806 <td class="paramname"><em>dst_stride_z</em>, </td>
807 </tr>
808 <tr>
809 <td class="paramkey"></td>
810 <td></td>
811 <td class="paramtype">uint&#160;</td>
812 <td class="paramname"><em>dst_step_z</em>, </td>
813 </tr>
814 <tr>
815 <td class="paramkey"></td>
816 <td></td>
817 <td class="paramtype">uint&#160;</td>
818 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
819 </tr>
820 <tr>
821 <td class="paramkey"></td>
822 <td></td>
823 <td class="paramtype">__global uchar *&#160;</td>
824 <td class="paramname"><em>sum_ptr</em>, </td>
825 </tr>
826 <tr>
827 <td class="paramkey"></td>
828 <td></td>
829 <td class="paramtype">uint&#160;</td>
830 <td class="paramname"><em>sum_stride_x</em>, </td>
831 </tr>
832 <tr>
833 <td class="paramkey"></td>
834 <td></td>
835 <td class="paramtype">uint&#160;</td>
836 <td class="paramname"><em>sum_step_x</em>, </td>
837 </tr>
838 <tr>
839 <td class="paramkey"></td>
840 <td></td>
841 <td class="paramtype">uint&#160;</td>
842 <td class="paramname"><em>sum_stride_y</em>, </td>
843 </tr>
844 <tr>
845 <td class="paramkey"></td>
846 <td></td>
847 <td class="paramtype">uint&#160;</td>
848 <td class="paramname"><em>sum_step_y</em>, </td>
849 </tr>
850 <tr>
851 <td class="paramkey"></td>
852 <td></td>
853 <td class="paramtype">uint&#160;</td>
854 <td class="paramname"><em>sum_stride_z</em>, </td>
855 </tr>
856 <tr>
857 <td class="paramkey"></td>
858 <td></td>
859 <td class="paramtype">uint&#160;</td>
860 <td class="paramname"><em>sum_step_z</em>, </td>
861 </tr>
862 <tr>
863 <td class="paramkey"></td>
864 <td></td>
865 <td class="paramtype">uint&#160;</td>
866 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
867 </tr>
868 <tr>
869 <td class="paramkey"></td>
870 <td></td>
871 <td class="paramtype">uint&#160;</td>
872 <td class="paramname"><em>width</em>&#160;</td>
873 </tr>
874 <tr>
875 <td></td>
876 <td>)</td>
877 <td></td><td></td>
878 </tr>
879 </table>
880</div><div class="memdoc">
881
882<p>Identifies the maximum value across the 1st dimension and shifts the values of the input tensor by this maximum value, then gets the exponent of each element as sums all elements across each row. </p>
883<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>
884<dd>
885Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
886<dd>
887In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
888<dd>
889Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
890<dl class="params"><dt>Parameters</dt><dd>
891 <table class="params">
892 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: QS8/QS16/F16/F32 </td></tr>
893 <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>
894 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_x</td><td>src_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
895 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source tensor in Y dimension (in bytes) </td></tr>
896 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_y</td><td>src_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
897 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
898 <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>
899 <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>
900 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_ptr</td><td>Pointer to the max values tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
901 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_stride_x</td><td>Stride of the max values tensor in X dimension (in bytes) </td></tr>
902 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_step_x</td><td>max_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
903 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_stride_y</td><td>Stride of the max values tensor in Y dimension (in bytes) </td></tr>
904 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_step_y</td><td>max_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
905 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_stride_z</td><td>Stride of the max values tensor in Z dimension (in bytes) </td></tr>
906 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_step_z</td><td>max_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
907 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_offset_first_element_in_bytes</td><td>The offset of the first element in the max values tensor </td></tr>
908 <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>
909 <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>
910 <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>
911 <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>
912 <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>
913 <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>
914 <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>
915 <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>
916 <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>
917 <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>
918 <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>
919 <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>
920 <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>
921 <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>
922 <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>
923 <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>
924 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
925 </table>
926 </dd>
927</dl>
928
929<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">500</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
930
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]931<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00112">arm_compute::test::validation::beta</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00123">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00048">EXP_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">GRID_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00088">idx4</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00046">MUL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">SELECT_DATA_TYPE</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00133">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00081">type_min_</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00040">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00043">VSTORE</a>.</p>
932<div class="fragment"><div class="line"><a name="l00506"></a><span class="lineno"> 506</span>&#160;{</div><div class="line"><a name="l00507"></a><span class="lineno"> 507</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00508"></a><span class="lineno"> 508</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00509"></a><span class="lineno"> 509</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> maxo = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(maxo);</div><div class="line"><a name="l00510"></a><span class="lineno"> 510</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><div class="line"><a name="l00511"></a><span class="lineno"> 511</span>&#160;</div><div class="line"><a name="l00512"></a><span class="lineno"> 512</span>&#160; <span class="keyword">const</span> uint lid = get_local_id(0);</div><div class="line"><a name="l00513"></a><span class="lineno"> 513</span>&#160;</div><div class="line"><a name="l00514"></a><span class="lineno"> 514</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00515"></a><span class="lineno"> 515</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00516"></a><span class="lineno"> 516</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>, 4)</div><div class="line"><a name="l00517"></a><span class="lineno"> 517</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a> = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4))BETA;</div><div class="line"><a name="l00518"></a><span class="lineno"> 518</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00519"></a><span class="lineno"> 519</span>&#160;</div><div class="line"><a name="l00520"></a><span class="lineno"> 520</span>&#160; <span class="comment">// Define one temporary vector per work-item.</span></div><div class="line"><a name="l00521"></a><span class="lineno"> 521</span>&#160; __local <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4) tmp_local[<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>];</div><div class="line"><a name="l00522"></a><span class="lineno"> 522</span>&#160; __local DATA_TYPE max_local;</div><div class="line"><a name="l00523"></a><span class="lineno"> 523</span>&#160;</div><div class="line"><a name="l00524"></a><span class="lineno"> 524</span>&#160; __constant <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4) type_min4 = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4))(<a class="code" href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a>);</div><div class="line"><a name="l00525"></a><span class="lineno"> 525</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00526"></a><span class="lineno"> 526</span>&#160; max_val_vec = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4))type_min4;</div><div class="line"><a name="l00527"></a><span class="lineno"> 527</span>&#160; <span class="comment">// Number of elements per work-item.</span></div><div class="line"><a name="l00528"></a><span class="lineno"> 528</span>&#160; const uint row = width / GRID_SIZE;</div><div class="line"><a name="l00529"></a><span class="lineno"> 529</span>&#160; <span class="comment">// Number of iterations per work-item.</span></div><div class="line"><a name="l00530"></a><span class="lineno"> 530</span>&#160; const uint width_ = row &gt;&gt; 2;</div><div class="line"><a name="l00531"></a><span class="lineno"> 531</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00532"></a><span class="lineno"> 532</span>&#160; uint i = 0;</div><div class="line"><a name="l00533"></a><span class="lineno"> 533</span>&#160; for(; i &lt; width_; i++)</div><div class="line"><a name="l00534"></a><span class="lineno"> 534</span>&#160; {</div><div class="line"><a name="l00535"></a><span class="lineno"> 535</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00536"></a><span class="lineno"> 536</span>&#160; data_max = <a class="code" href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(4)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i * GRID_SIZE * 4, 0));</div><div class="line"><a name="l00537"></a><span class="lineno"> 537</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data_max, max_val_vec, DATA_TYPE, 4);</div><div class="line"><a name="l00538"></a><span class="lineno"> 538</span>&#160; }</div><div class="line"><a name="l00539"></a><span class="lineno"> 539</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00540"></a><span class="lineno"> 540</span>&#160; <span class="comment">// How many work-items needed to complete the computation.</span></div><div class="line"><a name="l00541"></a><span class="lineno"> 541</span>&#160; <span class="keywordtype">int</span> boundary_workitems = (width % (GRID_SIZE * 4)) / 4;</div><div class="line"><a name="l00542"></a><span class="lineno"> 542</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00543"></a><span class="lineno"> 543</span>&#160; {</div><div class="line"><a name="l00544"></a><span class="lineno"> 544</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00545"></a><span class="lineno"> 545</span>&#160; data_max = VLOAD(4)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i * GRID_SIZE * 4, 0));</div><div class="line"><a name="l00546"></a><span class="lineno"> 546</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data_max, max_val_vec, DATA_TYPE, 4);</div><div class="line"><a name="l00547"></a><span class="lineno"> 547</span>&#160; }</div><div class="line"><a name="l00548"></a><span class="lineno"> 548</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00549"></a><span class="lineno"> 549</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00550"></a><span class="lineno"> 550</span>&#160; {</div><div class="line"><a name="l00551"></a><span class="lineno"> 551</span>&#160; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00552"></a><span class="lineno"> 552</span>&#160; i--;</div><div class="line"><a name="l00553"></a><span class="lineno"> 553</span>&#160; }</div><div class="line"><a name="l00554"></a><span class="lineno"> 554</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00555"></a><span class="lineno"> 555</span>&#160; {</div><div class="line"><a name="l00556"></a><span class="lineno"> 556</span>&#160; <span class="comment">// Handle non multiple of 4</span></div><div class="line"><a name="l00557"></a><span class="lineno"> 557</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00558"></a><span class="lineno"> 558</span>&#160; data_max = VLOAD(4)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, (GRID_SIZE * i * 4) + 4, 0));</div><div class="line"><a name="l00559"></a><span class="lineno"> 559</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>, 4)</div><div class="line"><a name="l00560"></a><span class="lineno"> 560</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint4)(GRID_SIZE * i * 4) + boundary_workitems * 4 + <a class="code" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4));</div><div class="line"><a name="l00561"></a><span class="lineno"> 561</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec, select(<a class="code" href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a>, data_max, widx), DATA_TYPE, 4);</div><div class="line"><a name="l00562"></a><span class="lineno"> 562</span>&#160; }</div><div class="line"><a name="l00563"></a><span class="lineno"> 563</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_VECTOR_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00564"></a><span class="lineno"> 564</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_GRID_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00565"></a><span class="lineno"> 565</span>&#160; tmp_local[lid] = max_val_vec;</div><div class="line"><a name="l00566"></a><span class="lineno"> 566</span>&#160;</div><div class="line"><a name="l00567"></a><span class="lineno"> 567</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00568"></a><span class="lineno"> 568</span>&#160;</div><div class="line"><a name="l00569"></a><span class="lineno"> 569</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 256)</div><div class="line"><a name="l00570"></a><span class="lineno"> 570</span>&#160; {</div><div class="line"><a name="l00571"></a><span class="lineno"> 571</span>&#160; <span class="keywordflow">if</span>(lid &lt; 128)</div><div class="line"><a name="l00572"></a><span class="lineno"> 572</span>&#160; {</div><div class="line"><a name="l00573"></a><span class="lineno"> 573</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 128], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00574"></a><span class="lineno"> 574</span>&#160; }</div><div class="line"><a name="l00575"></a><span class="lineno"> 575</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00576"></a><span class="lineno"> 576</span>&#160; }</div><div class="line"><a name="l00577"></a><span class="lineno"> 577</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 128)</div><div class="line"><a name="l00578"></a><span class="lineno"> 578</span>&#160; {</div><div class="line"><a name="l00579"></a><span class="lineno"> 579</span>&#160; <span class="keywordflow">if</span>(lid &lt; 64)</div><div class="line"><a name="l00580"></a><span class="lineno"> 580</span>&#160; {</div><div class="line"><a name="l00581"></a><span class="lineno"> 581</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 64], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00582"></a><span class="lineno"> 582</span>&#160; }</div><div class="line"><a name="l00583"></a><span class="lineno"> 583</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00584"></a><span class="lineno"> 584</span>&#160; }</div><div class="line"><a name="l00585"></a><span class="lineno"> 585</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 64)</div><div class="line"><a name="l00586"></a><span class="lineno"> 586</span>&#160; {</div><div class="line"><a name="l00587"></a><span class="lineno"> 587</span>&#160; <span class="keywordflow">if</span>(lid &lt; 32)</div><div class="line"><a name="l00588"></a><span class="lineno"> 588</span>&#160; {</div><div class="line"><a name="l00589"></a><span class="lineno"> 589</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 32], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00590"></a><span class="lineno"> 590</span>&#160; }</div><div class="line"><a name="l00591"></a><span class="lineno"> 591</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00592"></a><span class="lineno"> 592</span>&#160; }</div><div class="line"><a name="l00593"></a><span class="lineno"> 593</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 32)</div><div class="line"><a name="l00594"></a><span class="lineno"> 594</span>&#160; {</div><div class="line"><a name="l00595"></a><span class="lineno"> 595</span>&#160; <span class="keywordflow">if</span>(lid &lt; 16)</div><div class="line"><a name="l00596"></a><span class="lineno"> 596</span>&#160; {</div><div class="line"><a name="l00597"></a><span class="lineno"> 597</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 16], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00598"></a><span class="lineno"> 598</span>&#160; }</div><div class="line"><a name="l00599"></a><span class="lineno"> 599</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00600"></a><span class="lineno"> 600</span>&#160; }</div><div class="line"><a name="l00601"></a><span class="lineno"> 601</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 16)</div><div class="line"><a name="l00602"></a><span class="lineno"> 602</span>&#160; {</div><div class="line"><a name="l00603"></a><span class="lineno"> 603</span>&#160; <span class="keywordflow">if</span>(lid &lt; 8)</div><div class="line"><a name="l00604"></a><span class="lineno"> 604</span>&#160; {</div><div class="line"><a name="l00605"></a><span class="lineno"> 605</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 8], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00606"></a><span class="lineno"> 606</span>&#160; }</div><div class="line"><a name="l00607"></a><span class="lineno"> 607</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00608"></a><span class="lineno"> 608</span>&#160; }</div><div class="line"><a name="l00609"></a><span class="lineno"> 609</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 8)</div><div class="line"><a name="l00610"></a><span class="lineno"> 610</span>&#160; {</div><div class="line"><a name="l00611"></a><span class="lineno"> 611</span>&#160; <span class="keywordflow">if</span>(lid &lt; 4)</div><div class="line"><a name="l00612"></a><span class="lineno"> 612</span>&#160; {</div><div class="line"><a name="l00613"></a><span class="lineno"> 613</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 4], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00614"></a><span class="lineno"> 614</span>&#160; }</div><div class="line"><a name="l00615"></a><span class="lineno"> 615</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00616"></a><span class="lineno"> 616</span>&#160; }</div><div class="line"><a name="l00617"></a><span class="lineno"> 617</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 4)</div><div class="line"><a name="l00618"></a><span class="lineno"> 618</span>&#160; {</div><div class="line"><a name="l00619"></a><span class="lineno"> 619</span>&#160; <span class="keywordflow">if</span>(lid &lt; 2)</div><div class="line"><a name="l00620"></a><span class="lineno"> 620</span>&#160; {</div><div class="line"><a name="l00621"></a><span class="lineno"> 621</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 2], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00622"></a><span class="lineno"> 622</span>&#160; }</div><div class="line"><a name="l00623"></a><span class="lineno"> 623</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00624"></a><span class="lineno"> 624</span>&#160; }</div><div class="line"><a name="l00625"></a><span class="lineno"> 625</span>&#160; <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00626"></a><span class="lineno"> 626</span>&#160; {</div><div class="line"><a name="l00627"></a><span class="lineno"> 627</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 1], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00628"></a><span class="lineno"> 628</span>&#160; max_val_vec.s01 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec.s01, max_val_vec.s23, DATA_TYPE, 2);</div><div class="line"><a name="l00629"></a><span class="lineno"> 629</span>&#160; max_val_vec.s0 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec.s0, max_val_vec.s1, DATA_TYPE, 1);</div><div class="line"><a name="l00630"></a><span class="lineno"> 630</span>&#160; max_local = max_val_vec.s0;</div><div class="line"><a name="l00631"></a><span class="lineno"> 631</span>&#160; }</div><div class="line"><a name="l00632"></a><span class="lineno"> 632</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00633"></a><span class="lineno"> 633</span>&#160;</div><div class="line"><a name="l00634"></a><span class="lineno"> 634</span>&#160; <span class="comment">/* Second section */</span></div><div class="line"><a name="l00635"></a><span class="lineno"> 635</span>&#160;</div><div class="line"><a name="l00636"></a><span class="lineno"> 636</span>&#160; <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00637"></a><span class="lineno"> 637</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00638"></a><span class="lineno"> 638</span>&#160; sum1D = 0;</div><div class="line"><a name="l00639"></a><span class="lineno"> 639</span>&#160; DATA_TYPE max_val = max_local;</div><div class="line"><a name="l00640"></a><span class="lineno"> 640</span>&#160;</div><div class="line"><a name="l00641"></a><span class="lineno"> 641</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00642"></a><span class="lineno"> 642</span>&#160; for(i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00643"></a><span class="lineno"> 643</span>&#160; {</div><div class="line"><a name="l00644"></a><span class="lineno"> 644</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00645"></a><span class="lineno"> 645</span>&#160; data = VLOAD(4)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i * GRID_SIZE * 4, 0));</div><div class="line"><a name="l00646"></a><span class="lineno"> 646</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, 4);</div><div class="line"><a name="l00647"></a><span class="lineno"> 647</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00648"></a><span class="lineno"> 648</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a>, DATA_TYPE, 4);</div><div class="line"><a name="l00649"></a><span class="lineno"> 649</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00650"></a><span class="lineno"> 650</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, 4);</div><div class="line"><a name="l00651"></a><span class="lineno"> 651</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00652"></a><span class="lineno"> 652</span>&#160; (data, 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, i * GRID_SIZE * 4, 0));</div><div class="line"><a name="l00653"></a><span class="lineno"> 653</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, 4);</div><div class="line"><a name="l00654"></a><span class="lineno"> 654</span>&#160; }</div><div class="line"><a name="l00655"></a><span class="lineno"> 655</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00656"></a><span class="lineno"> 656</span>&#160; boundary_workitems = (width % (GRID_SIZE * 4)) / 4;</div><div class="line"><a name="l00657"></a><span class="lineno"> 657</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00658"></a><span class="lineno"> 658</span>&#160; {</div><div class="line"><a name="l00659"></a><span class="lineno"> 659</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00660"></a><span class="lineno"> 660</span>&#160; data = VLOAD(4)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i * GRID_SIZE * 4, 0));</div><div class="line"><a name="l00661"></a><span class="lineno"> 661</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, 4);</div><div class="line"><a name="l00662"></a><span class="lineno"> 662</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00663"></a><span class="lineno"> 663</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a>, DATA_TYPE, 4);</div><div class="line"><a name="l00664"></a><span class="lineno"> 664</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00665"></a><span class="lineno"> 665</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, 4);</div><div class="line"><a name="l00666"></a><span class="lineno"> 666</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00667"></a><span class="lineno"> 667</span>&#160; (data, 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, i * GRID_SIZE * 4, 0));</div><div class="line"><a name="l00668"></a><span class="lineno"> 668</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, 4);</div><div class="line"><a name="l00669"></a><span class="lineno"> 669</span>&#160; }</div><div class="line"><a name="l00670"></a><span class="lineno"> 670</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00671"></a><span class="lineno"> 671</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00672"></a><span class="lineno"> 672</span>&#160; {</div><div class="line"><a name="l00673"></a><span class="lineno"> 673</span>&#160; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00674"></a><span class="lineno"> 674</span>&#160; i--;</div><div class="line"><a name="l00675"></a><span class="lineno"> 675</span>&#160; }</div><div class="line"><a name="l00676"></a><span class="lineno"> 676</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00677"></a><span class="lineno"> 677</span>&#160; {</div><div class="line"><a name="l00678"></a><span class="lineno"> 678</span>&#160; <span class="comment">// Handle non multiple of vector size ((GRID_SIZE * i * 4) + 4, 0); move 4 float positions ahead, *4 is due to the stride</span></div><div class="line"><a name="l00679"></a><span class="lineno"> 679</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00680"></a><span class="lineno"> 680</span>&#160; data = VLOAD(4)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, (GRID_SIZE * i * 4) + 4, 0));</div><div class="line"><a name="l00681"></a><span class="lineno"> 681</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, 4);</div><div class="line"><a name="l00682"></a><span class="lineno"> 682</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00683"></a><span class="lineno"> 683</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a>, DATA_TYPE, 4);</div><div class="line"><a name="l00684"></a><span class="lineno"> 684</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00685"></a><span class="lineno"> 685</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, 4);</div><div class="line"><a name="l00686"></a><span class="lineno"> 686</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4)</div><div class="line"><a name="l00687"></a><span class="lineno"> 687</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint4)(GRID_SIZE * i * 4) + boundary_workitems * 4 + <a class="code" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4));</div><div class="line"><a name="l00688"></a><span class="lineno"> 688</span>&#160; data = select(0, data, widx);</div><div class="line"><a name="l00689"></a><span class="lineno"> 689</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00690"></a><span class="lineno"> 690</span>&#160; (data, 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, (GRID_SIZE * i * 4) + 4, 0));</div><div class="line"><a name="l00691"></a><span class="lineno"> 691</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, 4);</div><div class="line"><a name="l00692"></a><span class="lineno"> 692</span>&#160; }</div><div class="line"><a name="l00693"></a><span class="lineno"> 693</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_VECTOR_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00694"></a><span class="lineno"> 694</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_GRID_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00695"></a><span class="lineno"> 695</span>&#160; tmp_local[lid] = sum1D;</div><div class="line"><a name="l00696"></a><span class="lineno"> 696</span>&#160;</div><div class="line"><a name="l00697"></a><span class="lineno"> 697</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00698"></a><span class="lineno"> 698</span>&#160;</div><div class="line"><a name="l00699"></a><span class="lineno"> 699</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 256)</div><div class="line"><a name="l00700"></a><span class="lineno"> 700</span>&#160; {</div><div class="line"><a name="l00701"></a><span class="lineno"> 701</span>&#160; <span class="keywordflow">if</span>(lid &lt; 128)</div><div class="line"><a name="l00702"></a><span class="lineno"> 702</span>&#160; {</div><div class="line"><a name="l00703"></a><span class="lineno"> 703</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 128], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00704"></a><span class="lineno"> 704</span>&#160; }</div><div class="line"><a name="l00705"></a><span class="lineno"> 705</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00706"></a><span class="lineno"> 706</span>&#160; }</div><div class="line"><a name="l00707"></a><span class="lineno"> 707</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 128)</div><div class="line"><a name="l00708"></a><span class="lineno"> 708</span>&#160; {</div><div class="line"><a name="l00709"></a><span class="lineno"> 709</span>&#160; <span class="keywordflow">if</span>(lid &lt; 64)</div><div class="line"><a name="l00710"></a><span class="lineno"> 710</span>&#160; {</div><div class="line"><a name="l00711"></a><span class="lineno"> 711</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 64], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00712"></a><span class="lineno"> 712</span>&#160; }</div><div class="line"><a name="l00713"></a><span class="lineno"> 713</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00714"></a><span class="lineno"> 714</span>&#160; }</div><div class="line"><a name="l00715"></a><span class="lineno"> 715</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 64)</div><div class="line"><a name="l00716"></a><span class="lineno"> 716</span>&#160; {</div><div class="line"><a name="l00717"></a><span class="lineno"> 717</span>&#160; <span class="keywordflow">if</span>(lid &lt; 32)</div><div class="line"><a name="l00718"></a><span class="lineno"> 718</span>&#160; {</div><div class="line"><a name="l00719"></a><span class="lineno"> 719</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 32], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00720"></a><span class="lineno"> 720</span>&#160; }</div><div class="line"><a name="l00721"></a><span class="lineno"> 721</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00722"></a><span class="lineno"> 722</span>&#160; }</div><div class="line"><a name="l00723"></a><span class="lineno"> 723</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 32)</div><div class="line"><a name="l00724"></a><span class="lineno"> 724</span>&#160; {</div><div class="line"><a name="l00725"></a><span class="lineno"> 725</span>&#160; <span class="keywordflow">if</span>(lid &lt; 16)</div><div class="line"><a name="l00726"></a><span class="lineno"> 726</span>&#160; {</div><div class="line"><a name="l00727"></a><span class="lineno"> 727</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 16], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00728"></a><span class="lineno"> 728</span>&#160; }</div><div class="line"><a name="l00729"></a><span class="lineno"> 729</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00730"></a><span class="lineno"> 730</span>&#160; }</div><div class="line"><a name="l00731"></a><span class="lineno"> 731</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 16)</div><div class="line"><a name="l00732"></a><span class="lineno"> 732</span>&#160; {</div><div class="line"><a name="l00733"></a><span class="lineno"> 733</span>&#160; <span class="keywordflow">if</span>(lid &lt; 8)</div><div class="line"><a name="l00734"></a><span class="lineno"> 734</span>&#160; {</div><div class="line"><a name="l00735"></a><span class="lineno"> 735</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 8], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00736"></a><span class="lineno"> 736</span>&#160; }</div><div class="line"><a name="l00737"></a><span class="lineno"> 737</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00738"></a><span class="lineno"> 738</span>&#160; }</div><div class="line"><a name="l00739"></a><span class="lineno"> 739</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 8)</div><div class="line"><a name="l00740"></a><span class="lineno"> 740</span>&#160; {</div><div class="line"><a name="l00741"></a><span class="lineno"> 741</span>&#160; <span class="keywordflow">if</span>(lid &lt; 4)</div><div class="line"><a name="l00742"></a><span class="lineno"> 742</span>&#160; {</div><div class="line"><a name="l00743"></a><span class="lineno"> 743</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 4], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00744"></a><span class="lineno"> 744</span>&#160; }</div><div class="line"><a name="l00745"></a><span class="lineno"> 745</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00746"></a><span class="lineno"> 746</span>&#160; }</div><div class="line"><a name="l00747"></a><span class="lineno"> 747</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 4)</div><div class="line"><a name="l00748"></a><span class="lineno"> 748</span>&#160; {</div><div class="line"><a name="l00749"></a><span class="lineno"> 749</span>&#160; <span class="keywordflow">if</span>(lid &lt; 2)</div><div class="line"><a name="l00750"></a><span class="lineno"> 750</span>&#160; {</div><div class="line"><a name="l00751"></a><span class="lineno"> 751</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 2], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00752"></a><span class="lineno"> 752</span>&#160; }</div><div class="line"><a name="l00753"></a><span class="lineno"> 753</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00754"></a><span class="lineno"> 754</span>&#160; }</div><div class="line"><a name="l00755"></a><span class="lineno"> 755</span>&#160; <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00756"></a><span class="lineno"> 756</span>&#160; {</div><div class="line"><a name="l00757"></a><span class="lineno"> 757</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 1], tmp_local[lid], DATA_TYPE, 4);</div><div class="line"><a name="l00758"></a><span class="lineno"> 758</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00759"></a><span class="lineno"> 759</span>&#160; sum1D.s01 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01, sum1D.s23, DATA_TYPE, 2);</div><div class="line"><a name="l00760"></a><span class="lineno"> 760</span>&#160; sum1D.s0 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0, sum1D.s1, DATA_TYPE, 1);</div><div class="line"><a name="l00761"></a><span class="lineno"> 761</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div><div class="line"><a name="l00762"></a><span class="lineno"> 762</span>&#160; }</div><div class="line"><a name="l00763"></a><span class="lineno"> 763</span>&#160;}</div><div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ab146b9cbab6e73e7588b240dc709fe01"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">arm_compute::test::validation::beta</a></div><div class="ttdeci">beta</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00112">GEMM.cpp:112</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]933<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>
934<div class="ttc" id="helpers_8h_xhtml_a287e2fc366c312b468382c95bb90f91f"><div class="ttname"><a href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a></div><div class="ttdeci">#define VLOAD(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00040">helpers.h:40</a></div></div>
935<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#l00048">softmax_layer.cl:48</a></div></div>
936<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>
937<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#l00043">softmax_layer.cl:43</a></div></div>
938<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>
939<div class="ttc" id="softmax__layer_8cl_xhtml_af7a08044d0e491a0ee1520a24a107a2b"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">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#l00081">softmax_layer.cl:81</a></div></div>
940<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#l00044">softmax_layer.cl:44</a></div></div>
941<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
942<div class="ttc" id="helpers_8h_xhtml_acb282042d1edeeaa3cc979a206f78b54"><div class="ttname"><a href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a></div><div class="ttdeci">#define VSTORE(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00043">helpers.h:43</a></div></div>
943<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#l00055">softmax_layer.cl:55</a></div></div>
944<div class="ttc" id="softmax__layer_8cl_xhtml_a4884a666a1e93fbf8c27bd7d2da3c8bb"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a></div><div class="ttdeci">__constant uint4 idx4</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00088">softmax_layer.cl:88</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]945<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00123">Convolution.cpp:123</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]946<div class="ttc" id="softmax__layer_8cl_xhtml_a22303c4047ec5027c1538d53964b9d0d"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00046">softmax_layer.cl:46</a></div></div>
947<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
948<div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00142">helpers.h:142</a></div></div>
949<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#l00045">softmax_layer.cl:45</a></div></div>
950<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>
951<div class="ttc" id="softmax__layer_8cl_xhtml_a80b22c555ddadb47cc6ca338a9c49126"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a></div><div class="ttdeci">#define MINVAL</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00054">softmax_layer.cl:54</a></div></div>
952<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
953<div class="ttc" id="softmax__layer_8cl_xhtml_a08246606c233e7785a497c09672f366f"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a></div><div class="ttdeci">#define GRID_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00062">softmax_layer.cl:62</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]954<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00133">Convolution.cpp:133</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]955</div><!-- fragment -->
956</div>
957</div>
958<a class="anchor" id="a9d74601bcabbb2f14bcf52385ad666dc"></a>
959<div class="memitem">
960<div class="memproto">
961 <table class="memname">
962 <tr>
963 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_serial </td>
964 <td>(</td>
965 <td class="paramtype">__global uchar *&#160;</td>
966 <td class="paramname"><em>src_ptr</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>src_stride_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>src_step_x</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>src_stride_y</em>, </td>
985 </tr>
986 <tr>
987 <td class="paramkey"></td>
988 <td></td>
989 <td class="paramtype">uint&#160;</td>
990 <td class="paramname"><em>src_step_y</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>src_stride_z</em>, </td>
997 </tr>
998 <tr>
999 <td class="paramkey"></td>
1000 <td></td>
1001 <td class="paramtype">uint&#160;</td>
1002 <td class="paramname"><em>src_step_z</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>src_offset_first_element_in_bytes</em>, </td>
1009 </tr>
1010 <tr>
1011 <td class="paramkey"></td>
1012 <td></td>
1013 <td class="paramtype">__global uchar *&#160;</td>
1014 <td class="paramname"><em>maxo_ptr</em>, </td>
1015 </tr>
1016 <tr>
1017 <td class="paramkey"></td>
1018 <td></td>
1019 <td class="paramtype">uint&#160;</td>
1020 <td class="paramname"><em>maxo_stride_x</em>, </td>
1021 </tr>
1022 <tr>
1023 <td class="paramkey"></td>
1024 <td></td>
1025 <td class="paramtype">uint&#160;</td>
1026 <td class="paramname"><em>maxo_step_x</em>, </td>
1027 </tr>
1028 <tr>
1029 <td class="paramkey"></td>
1030 <td></td>
1031 <td class="paramtype">uint&#160;</td>
1032 <td class="paramname"><em>maxo_stride_y</em>, </td>
1033 </tr>
1034 <tr>
1035 <td class="paramkey"></td>
1036 <td></td>
1037 <td class="paramtype">uint&#160;</td>
1038 <td class="paramname"><em>maxo_step_y</em>, </td>
1039 </tr>
1040 <tr>
1041 <td class="paramkey"></td>
1042 <td></td>
1043 <td class="paramtype">uint&#160;</td>
1044 <td class="paramname"><em>maxo_stride_z</em>, </td>
1045 </tr>
1046 <tr>
1047 <td class="paramkey"></td>
1048 <td></td>
1049 <td class="paramtype">uint&#160;</td>
1050 <td class="paramname"><em>maxo_step_z</em>, </td>
1051 </tr>
1052 <tr>
1053 <td class="paramkey"></td>
1054 <td></td>
1055 <td class="paramtype">uint&#160;</td>
1056 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
1057 </tr>
1058 <tr>
1059 <td class="paramkey"></td>
1060 <td></td>
1061 <td class="paramtype">__global uchar *&#160;</td>
1062 <td class="paramname"><em>dst_ptr</em>, </td>
1063 </tr>
1064 <tr>
1065 <td class="paramkey"></td>
1066 <td></td>
1067 <td class="paramtype">uint&#160;</td>
1068 <td class="paramname"><em>dst_stride_x</em>, </td>
1069 </tr>
1070 <tr>
1071 <td class="paramkey"></td>
1072 <td></td>
1073 <td class="paramtype">uint&#160;</td>
1074 <td class="paramname"><em>dst_step_x</em>, </td>
1075 </tr>
1076 <tr>
1077 <td class="paramkey"></td>
1078 <td></td>
1079 <td class="paramtype">uint&#160;</td>
1080 <td class="paramname"><em>dst_stride_y</em>, </td>
1081 </tr>
1082 <tr>
1083 <td class="paramkey"></td>
1084 <td></td>
1085 <td class="paramtype">uint&#160;</td>
1086 <td class="paramname"><em>dst_step_y</em>, </td>
1087 </tr>
1088 <tr>
1089 <td class="paramkey"></td>
1090 <td></td>
1091 <td class="paramtype">uint&#160;</td>
1092 <td class="paramname"><em>dst_stride_z</em>, </td>
1093 </tr>
1094 <tr>
1095 <td class="paramkey"></td>
1096 <td></td>
1097 <td class="paramtype">uint&#160;</td>
1098 <td class="paramname"><em>dst_step_z</em>, </td>
1099 </tr>
1100 <tr>
1101 <td class="paramkey"></td>
1102 <td></td>
1103 <td class="paramtype">uint&#160;</td>
1104 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
1105 </tr>
1106 <tr>
1107 <td class="paramkey"></td>
1108 <td></td>
1109 <td class="paramtype">__global uchar *&#160;</td>
1110 <td class="paramname"><em>sum_ptr</em>, </td>
1111 </tr>
1112 <tr>
1113 <td class="paramkey"></td>
1114 <td></td>
1115 <td class="paramtype">uint&#160;</td>
1116 <td class="paramname"><em>sum_stride_x</em>, </td>
1117 </tr>
1118 <tr>
1119 <td class="paramkey"></td>
1120 <td></td>
1121 <td class="paramtype">uint&#160;</td>
1122 <td class="paramname"><em>sum_step_x</em>, </td>
1123 </tr>
1124 <tr>
1125 <td class="paramkey"></td>
1126 <td></td>
1127 <td class="paramtype">uint&#160;</td>
1128 <td class="paramname"><em>sum_stride_y</em>, </td>
1129 </tr>
1130 <tr>
1131 <td class="paramkey"></td>
1132 <td></td>
1133 <td class="paramtype">uint&#160;</td>
1134 <td class="paramname"><em>sum_step_y</em>, </td>
1135 </tr>
1136 <tr>
1137 <td class="paramkey"></td>
1138 <td></td>
1139 <td class="paramtype">uint&#160;</td>
1140 <td class="paramname"><em>sum_stride_z</em>, </td>
1141 </tr>
1142 <tr>
1143 <td class="paramkey"></td>
1144 <td></td>
1145 <td class="paramtype">uint&#160;</td>
1146 <td class="paramname"><em>sum_step_z</em>, </td>
1147 </tr>
1148 <tr>
1149 <td class="paramkey"></td>
1150 <td></td>
1151 <td class="paramtype">uint&#160;</td>
1152 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1153 </tr>
1154 <tr>
1155 <td class="paramkey"></td>
1156 <td></td>
1157 <td class="paramtype">uint&#160;</td>
1158 <td class="paramname"><em>width</em>&#160;</td>
1159 </tr>
1160 <tr>
1161 <td></td>
1162 <td>)</td>
1163 <td></td><td></td>
1164 </tr>
1165 </table>
1166</div><div class="memdoc">
1167
1168<p>Identifies the maximum value across the 1st dimension and shifts the values of the input tensor by this maximum value, then gets the exponent of each element as sums all elements across each row. </p>
1169<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>
1170<dd>
1171Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
1172<dd>
1173In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
1174<dd>
1175Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
1176<dl class="params"><dt>Parameters</dt><dd>
1177 <table class="params">
1178 <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>
1179 <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>
1180 <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>
1181 <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>
1182 <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>
1183 <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>
1184 <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>
1185 <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>
1186 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_ptr</td><td>Pointer to the max values tensor slice. Supported data types: same as <code>src_ptr</code> </td></tr>
1187 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_stride_x</td><td>Stride of the max values tensor in X dimension (in bytes) </td></tr>
1188 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_step_x</td><td>max_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
1189 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_stride_y</td><td>Stride of the max values tensor in Y dimension (in bytes) </td></tr>
1190 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_step_y</td><td>max_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
1191 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_stride_z</td><td>Stride of the max values tensor in Z dimension (in bytes) </td></tr>
1192 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_step_z</td><td>max_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
1193 <tr><td class="paramdir">[in]</td><td class="paramname">maxo_offset_first_element_in_bytes</td><td>The offset of the first element in the max values tensor </td></tr>
1194 <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>
1195 <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>
1196 <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>
1197 <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>
1198 <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>
1199 <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>
1200 <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>
1201 <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>
1202 <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>
1203 <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>
1204 <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>
1205 <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>
1206 <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>
1207 <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>
1208 <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>
1209 <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>
1210 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1211 </table>
1212 </dd>
1213</dl>
1214
1215<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">350</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1216
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1217<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00112">arm_compute::test::validation::beta</a>, <a class="el" href="helpers_8h_source.xhtml#l00049">CL_VEC_DATA_TYPE</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00123">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00048">EXP_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00035">EXPAND</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00082">idx__</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00080">LOG_VECTOR_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00046">MUL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">SELECT_DATA_TYPE</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00133">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00081">type_min_</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00079">VECTOR_SIZE</a>, <a class="el" href="helpers_8h_source.xhtml#l00040">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00043">VSTORE</a>.</p>
1218<div class="fragment"><div class="line"><a name="l00356"></a><span class="lineno"> 356</span>&#160;{</div><div class="line"><a name="l00357"></a><span class="lineno"> 357</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> maxo = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(maxo);</div><div class="line"><a name="l00360"></a><span class="lineno"> 360</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><div class="line"><a name="l00361"></a><span class="lineno"> 361</span>&#160;</div><div class="line"><a name="l00362"></a><span class="lineno"> 362</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00363"></a><span class="lineno"> 363</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00364"></a><span class="lineno"> 364</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>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00365"></a><span class="lineno"> 365</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a> = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>))BETA;</div><div class="line"><a name="l00366"></a><span class="lineno"> 366</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00367"></a><span class="lineno"> 367</span>&#160;</div><div class="line"><a name="l00368"></a><span class="lineno"> 368</span>&#160; <span class="comment">// Initialize local maximum</span></div><div class="line"><a name="l00369"></a><span class="lineno"> 369</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span>&#160; max_val_vec = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>))<a class="code" href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a>;</div><div class="line"><a name="l00371"></a><span class="lineno"> 371</span>&#160;</div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span>&#160; const uint width_ = width &gt;&gt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>;</div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span>&#160; for(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>&#160; {</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>&#160; data_max = <a class="code" href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; LOG_VECTOR_SIZE, 0));</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data_max, max_val_vec, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span>&#160; }</div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>&#160;</div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>&#160; data_max = <a class="code" href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(VECTOR_SIZE)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width_ &lt;&lt; LOG_VECTOR_SIZE, 0));</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</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>, VECTOR_SIZE)</div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((<a class="code" href="helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a>((<a class="code" href="helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a>(uint, VECTOR_SIZE)))(width_ &lt;&lt; LOG_VECTOR_SIZE) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, VECTOR_SIZE));</div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec, select(type_min_, data_max, widx), DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_VECTOR_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>&#160;</div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00390"></a><span class="lineno"> 390</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00391"></a><span class="lineno"> 391</span>&#160; max_val_vec.s01234567 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec.s01234567, max_val_vec.s89ABCDEF, DATA_TYPE, 8);</div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* VECTOR SIZE 16 END */</span><span class="preprocessor"></span></div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span>&#160; max_val_vec.s0123 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec.s0123, max_val_vec.s4567, DATA_TYPE, 4);</div><div class="line"><a name="l00395"></a><span class="lineno"> 395</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* VECTOR SIZE 8 END */</span><span class="preprocessor"></span></div><div class="line"><a name="l00396"></a><span class="lineno"> 396</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span>&#160; max_val_vec.s01 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec.s01, max_val_vec.s23, DATA_TYPE, 2);</div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* VECTOR SIZE 4 END */</span><span class="preprocessor"></span></div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>&#160; max_val_vec.s0 = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec.s0, max_val_vec.s1, DATA_TYPE, 1);</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span>&#160; <span class="comment">// Store result</span></div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>&#160; *((__global DATA_TYPE *)maxo.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = max_val_vec.s0;</div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span>&#160;</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>&#160; <span class="comment">/* Second section */</span></div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>&#160;</div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00406"></a><span class="lineno"> 406</span>&#160; DATA_TYPE max_val = *((__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;maxo, 0, 0));</div><div class="line"><a name="l00407"></a><span class="lineno"> 407</span>&#160;</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span>&#160; <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span>&#160; sum1D = 0;</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span>&#160;</div><div class="line"><a name="l00412"></a><span class="lineno"> 412</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>&#160; for(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00414"></a><span class="lineno"> 414</span>&#160; {</div><div class="line"><a name="l00415"></a><span class="lineno"> 415</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00416"></a><span class="lineno"> 416</span>&#160; data = <a class="code" href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(VECTOR_SIZE)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; LOG_VECTOR_SIZE, 0));</div><div class="line"><a name="l00417"></a><span class="lineno"> 417</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00418"></a><span class="lineno"> 418</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00419"></a><span class="lineno"> 419</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a>, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00420"></a><span class="lineno"> 420</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00421"></a><span class="lineno"> 421</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00422"></a><span class="lineno"> 422</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VECTOR_SIZE)</div><div class="line"><a name="l00423"></a><span class="lineno"> 423</span>&#160; (data, 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, i &lt;&lt; LOG_VECTOR_SIZE, 0));</div><div class="line"><a name="l00424"></a><span class="lineno"> 424</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00425"></a><span class="lineno"> 425</span>&#160; }</div><div class="line"><a name="l00426"></a><span class="lineno"> 426</span>&#160;</div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00428"></a><span class="lineno"> 428</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00429"></a><span class="lineno"> 429</span>&#160; data = <a class="code" href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(VECTOR_SIZE)(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width_ &lt;&lt; LOG_VECTOR_SIZE, 0));</div><div class="line"><a name="l00430"></a><span class="lineno"> 430</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00431"></a><span class="lineno"> 431</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00432"></a><span class="lineno"> 432</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a>, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00433"></a><span class="lineno"> 433</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00434"></a><span class="lineno"> 434</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00435"></a><span class="lineno"> 435</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((<a class="code" href="helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a>((<a class="code" href="helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a>(uint, VECTOR_SIZE)))(width_ &lt;&lt; LOG_VECTOR_SIZE) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, VECTOR_SIZE));</div><div class="line"><a name="l00436"></a><span class="lineno"> 436</span>&#160; data = select(0, data, widx);</div><div class="line"><a name="l00437"></a><span class="lineno"> 437</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VECTOR_SIZE)</div><div class="line"><a name="l00438"></a><span class="lineno"> 438</span>&#160; (data, 0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;dst, width_ &lt;&lt; LOG_VECTOR_SIZE, 0));</div><div class="line"><a name="l00439"></a><span class="lineno"> 439</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, VECTOR_SIZE);</div><div class="line"><a name="l00440"></a><span class="lineno"> 440</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_VECTOR_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00441"></a><span class="lineno"> 441</span>&#160;</div><div class="line"><a name="l00442"></a><span class="lineno"> 442</span>&#160; <span class="comment">// Perform sum reduction</span></div><div class="line"><a name="l00443"></a><span class="lineno"> 443</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00444"></a><span class="lineno"> 444</span>&#160; sum1D.s01234567 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01234567, sum1D.s89ABCDEF, DATA_TYPE, 8);</div><div class="line"><a name="l00445"></a><span class="lineno"> 445</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* VECTOR SIZE 16 END */</span><span class="preprocessor"></span></div><div class="line"><a name="l00446"></a><span class="lineno"> 446</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></div><div class="line"><a name="l00447"></a><span class="lineno"> 447</span>&#160; sum1D.s0123 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0123, sum1D.s4567, DATA_TYPE, 4);</div><div class="line"><a name="l00448"></a><span class="lineno"> 448</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* VECTOR SIZE 8 END */</span><span class="preprocessor"></span></div><div class="line"><a name="l00449"></a><span class="lineno"> 449</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00450"></a><span class="lineno"> 450</span>&#160; sum1D.s01 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01, sum1D.s23, DATA_TYPE, 2);</div><div class="line"><a name="l00451"></a><span class="lineno"> 451</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* VECTOR SIZE 4 END */</span><span class="preprocessor"></span></div><div class="line"><a name="l00452"></a><span class="lineno"> 452</span>&#160; sum1D.s0 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0, sum1D.s1, DATA_TYPE, 1);</div><div class="line"><a name="l00453"></a><span class="lineno"> 453</span>&#160;</div><div class="line"><a name="l00454"></a><span class="lineno"> 454</span>&#160; <span class="comment">// Calculate and store result</span></div><div class="line"><a name="l00455"></a><span class="lineno"> 455</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div><div class="line"><a name="l00456"></a><span class="lineno"> 456</span>&#160;}</div><div class="ttc" id="softmax__layer_8cl_xhtml_aa1dd94b8d98f1c6d790bdf0fc5de29e9"><div class="ttname"><a href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a></div><div class="ttdeci">__constant uint16 idx__</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00082">softmax_layer.cl:82</a></div></div>
1219<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ab146b9cbab6e73e7588b240dc709fe01"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">arm_compute::test::validation::beta</a></div><div class="ttdeci">beta</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00112">GEMM.cpp:112</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1220<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>
1221<div class="ttc" id="helpers_8h_xhtml_a287e2fc366c312b468382c95bb90f91f"><div class="ttname"><a href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a></div><div class="ttdeci">#define VLOAD(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00040">helpers.h:40</a></div></div>
1222<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#l00048">softmax_layer.cl:48</a></div></div>
1223<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>
1224<div class="ttc" id="softmax__layer_8cl_xhtml_a372393c380805985b813dbb16d589a64"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a></div><div class="ttdeci">#define LOG_VECTOR_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00080">softmax_layer.cl:80</a></div></div>
1225<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#l00043">softmax_layer.cl:43</a></div></div>
1226<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>
1227<div class="ttc" id="softmax__layer_8cl_xhtml_af7a08044d0e491a0ee1520a24a107a2b"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">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#l00081">softmax_layer.cl:81</a></div></div>
1228<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#l00044">softmax_layer.cl:44</a></div></div>
1229<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
1230<div class="ttc" id="helpers_8h_xhtml_acb282042d1edeeaa3cc979a206f78b54"><div class="ttname"><a href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a></div><div class="ttdeci">#define VSTORE(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00043">helpers.h:43</a></div></div>
1231<div class="ttc" id="softmax__layer_8cl_xhtml_a7c78836761fa3b5b124efea237dac70f"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a></div><div class="ttdeci">#define VECTOR_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00079">softmax_layer.cl:79</a></div></div>
1232<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#l00055">softmax_layer.cl:55</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1233<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00123">Convolution.cpp:123</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1234<div class="ttc" id="softmax__layer_8cl_xhtml_a22303c4047ec5027c1538d53964b9d0d"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00046">softmax_layer.cl:46</a></div></div>
1235<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
1236<div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00142">helpers.h:142</a></div></div>
1237<div class="ttc" id="helpers_8h_xhtml_a0f6ad555d097377151003a5119ccee45"><div class="ttname"><a href="helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a></div><div class="ttdeci">#define CL_VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00049">helpers.h:49</a></div></div>
1238<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#l00045">softmax_layer.cl:45</a></div></div>
1239<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>
1240<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
1241<div class="ttc" id="helpers_8h_xhtml_ae4b532a93c757194ec73b6790a3e6b1f"><div class="ttname"><a href="helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a></div><div class="ttdeci">#define EXPAND(x)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00035">helpers.h:35</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1242<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00133">Convolution.cpp:133</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1243</div><!-- fragment -->
1244</div>
1245</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1246<a class="anchor" id="ac4247ac0991e85965b7ded764e78f12c"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1247<div class="memitem">
1248<div class="memproto">
1249 <table class="memname">
1250 <tr>
1251 <td class="memname">__kernel void softmax_layer_norm </td>
1252 <td>(</td>
1253 <td class="paramtype">__global uchar *&#160;</td>
1254 <td class="paramname"><em>src_ptr</em>, </td>
1255 </tr>
1256 <tr>
1257 <td class="paramkey"></td>
1258 <td></td>
1259 <td class="paramtype">uint&#160;</td>
1260 <td class="paramname"><em>src_stride_x</em>, </td>
1261 </tr>
1262 <tr>
1263 <td class="paramkey"></td>
1264 <td></td>
1265 <td class="paramtype">uint&#160;</td>
1266 <td class="paramname"><em>src_step_x</em>, </td>
1267 </tr>
1268 <tr>
1269 <td class="paramkey"></td>
1270 <td></td>
1271 <td class="paramtype">uint&#160;</td>
1272 <td class="paramname"><em>src_stride_y</em>, </td>
1273 </tr>
1274 <tr>
1275 <td class="paramkey"></td>
1276 <td></td>
1277 <td class="paramtype">uint&#160;</td>
1278 <td class="paramname"><em>src_step_y</em>, </td>
1279 </tr>
1280 <tr>
1281 <td class="paramkey"></td>
1282 <td></td>
1283 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1284 <td class="paramname"><em>src_stride_z</em>, </td>
1285 </tr>
1286 <tr>
1287 <td class="paramkey"></td>
1288 <td></td>
1289 <td class="paramtype">uint&#160;</td>
1290 <td class="paramname"><em>src_step_z</em>, </td>
1291 </tr>
1292 <tr>
1293 <td class="paramkey"></td>
1294 <td></td>
1295 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1296 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1297 </tr>
1298 <tr>
1299 <td class="paramkey"></td>
1300 <td></td>
1301 <td class="paramtype">__global uchar *&#160;</td>
1302 <td class="paramname"><em>sum_ptr</em>, </td>
1303 </tr>
1304 <tr>
1305 <td class="paramkey"></td>
1306 <td></td>
1307 <td class="paramtype">uint&#160;</td>
1308 <td class="paramname"><em>sum_stride_x</em>, </td>
1309 </tr>
1310 <tr>
1311 <td class="paramkey"></td>
1312 <td></td>
1313 <td class="paramtype">uint&#160;</td>
1314 <td class="paramname"><em>sum_step_x</em>, </td>
1315 </tr>
1316 <tr>
1317 <td class="paramkey"></td>
1318 <td></td>
1319 <td class="paramtype">uint&#160;</td>
1320 <td class="paramname"><em>sum_stride_y</em>, </td>
1321 </tr>
1322 <tr>
1323 <td class="paramkey"></td>
1324 <td></td>
1325 <td class="paramtype">uint&#160;</td>
1326 <td class="paramname"><em>sum_step_y</em>, </td>
1327 </tr>
1328 <tr>
1329 <td class="paramkey"></td>
1330 <td></td>
1331 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1332 <td class="paramname"><em>sum_stride_z</em>, </td>
1333 </tr>
1334 <tr>
1335 <td class="paramkey"></td>
1336 <td></td>
1337 <td class="paramtype">uint&#160;</td>
1338 <td class="paramname"><em>sum_step_z</em>, </td>
1339 </tr>
1340 <tr>
1341 <td class="paramkey"></td>
1342 <td></td>
1343 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1344 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1345 </tr>
1346 <tr>
1347 <td class="paramkey"></td>
1348 <td></td>
1349 <td class="paramtype">__global uchar *&#160;</td>
1350 <td class="paramname"><em>dst_ptr</em>, </td>
1351 </tr>
1352 <tr>
1353 <td class="paramkey"></td>
1354 <td></td>
1355 <td class="paramtype">uint&#160;</td>
1356 <td class="paramname"><em>dst_stride_x</em>, </td>
1357 </tr>
1358 <tr>
1359 <td class="paramkey"></td>
1360 <td></td>
1361 <td class="paramtype">uint&#160;</td>
1362 <td class="paramname"><em>dst_step_x</em>, </td>
1363 </tr>
1364 <tr>
1365 <td class="paramkey"></td>
1366 <td></td>
1367 <td class="paramtype">uint&#160;</td>
1368 <td class="paramname"><em>dst_stride_y</em>, </td>
1369 </tr>
1370 <tr>
1371 <td class="paramkey"></td>
1372 <td></td>
1373 <td class="paramtype">uint&#160;</td>
1374 <td class="paramname"><em>dst_step_y</em>, </td>
1375 </tr>
1376 <tr>
1377 <td class="paramkey"></td>
1378 <td></td>
1379 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1380 <td class="paramname"><em>dst_stride_z</em>, </td>
1381 </tr>
1382 <tr>
1383 <td class="paramkey"></td>
1384 <td></td>
1385 <td class="paramtype">uint&#160;</td>
1386 <td class="paramname"><em>dst_step_z</em>, </td>
1387 </tr>
1388 <tr>
1389 <td class="paramkey"></td>
1390 <td></td>
1391 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1392 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
1393 </tr>
1394 <tr>
1395 <td></td>
1396 <td>)</td>
1397 <td></td><td></td>
1398 </tr>
1399 </table>
1400</div><div class="memdoc">
1401
1402<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]1403<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>
1404<dd>
1405Fixed 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]1406<dl class="params"><dt>Parameters</dt><dd>
1407 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1408 <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]1409 <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>
1410 <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>
1411 <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>
1412 <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]1413 <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>
1414 <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]1415 <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]1416 <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]1417 <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>
1418 <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>
1419 <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>
1420 <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]1421 <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>
1422 <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]1423 <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]1424 <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]1425 <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>
1426 <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>
1427 <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>
1428 <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]1429 <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>
1430 <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]1431 <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>
1432 </table>
1433 </dd>
1434</dl>
1435
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1436<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00292">292</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]1437
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1438<p>References <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00111">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00047">DIV_OP</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00123">arm_compute::test::validation::dst</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00133">arm_compute::test::validation::src</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>
1439<div class="fragment"><div class="line"><a name="l00296"></a><span class="lineno"> 296</span>&#160;{</div><div class="line"><a name="l00297"></a><span class="lineno"> 297</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00298"></a><span class="lineno"> 298</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00299"></a><span class="lineno"> 299</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><div class="line"><a name="l00300"></a><span class="lineno"> 300</span>&#160;</div><div class="line"><a name="l00301"></a><span class="lineno"> 301</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00302"></a><span class="lineno"> 302</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><div class="line"><a name="l00303"></a><span class="lineno"> 303</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><div class="line"><a name="l00304"></a><span class="lineno"> 304</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><div class="line"><a name="l00305"></a><span class="lineno"> 305</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><div class="line"><a name="l00306"></a><span class="lineno"> 306</span>&#160;}</div><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#l00111">helpers.h:111</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1440<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>
1441<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1442<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#l00047">softmax_layer.cl:47</a></div></div>
1443<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1444<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00123">Convolution.cpp:123</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1445<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1446<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1447<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1448<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00133">Convolution.cpp:133</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1449</div><!-- fragment -->
1450</div>
1451</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1452<a class="anchor" id="aa4412bb319c67deaeab07a481d93dfdb"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1453<div class="memitem">
1454<div class="memproto">
1455 <table class="memname">
1456 <tr>
1457 <td class="memname">__kernel void softmax_layer_shift_exp_sum </td>
1458 <td>(</td>
1459 <td class="paramtype">__global uchar *&#160;</td>
1460 <td class="paramname"><em>src_ptr</em>, </td>
1461 </tr>
1462 <tr>
1463 <td class="paramkey"></td>
1464 <td></td>
1465 <td class="paramtype">uint&#160;</td>
1466 <td class="paramname"><em>src_stride_x</em>, </td>
1467 </tr>
1468 <tr>
1469 <td class="paramkey"></td>
1470 <td></td>
1471 <td class="paramtype">uint&#160;</td>
1472 <td class="paramname"><em>src_step_x</em>, </td>
1473 </tr>
1474 <tr>
1475 <td class="paramkey"></td>
1476 <td></td>
1477 <td class="paramtype">uint&#160;</td>
1478 <td class="paramname"><em>src_stride_y</em>, </td>
1479 </tr>
1480 <tr>
1481 <td class="paramkey"></td>
1482 <td></td>
1483 <td class="paramtype">uint&#160;</td>
1484 <td class="paramname"><em>src_step_y</em>, </td>
1485 </tr>
1486 <tr>
1487 <td class="paramkey"></td>
1488 <td></td>
1489 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1490 <td class="paramname"><em>src_stride_z</em>, </td>
1491 </tr>
1492 <tr>
1493 <td class="paramkey"></td>
1494 <td></td>
1495 <td class="paramtype">uint&#160;</td>
1496 <td class="paramname"><em>src_step_z</em>, </td>
1497 </tr>
1498 <tr>
1499 <td class="paramkey"></td>
1500 <td></td>
1501 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1502 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1503 </tr>
1504 <tr>
1505 <td class="paramkey"></td>
1506 <td></td>
1507 <td class="paramtype">__global uchar *&#160;</td>
1508 <td class="paramname"><em>max_ptr</em>, </td>
1509 </tr>
1510 <tr>
1511 <td class="paramkey"></td>
1512 <td></td>
1513 <td class="paramtype">uint&#160;</td>
1514 <td class="paramname"><em>max_stride_x</em>, </td>
1515 </tr>
1516 <tr>
1517 <td class="paramkey"></td>
1518 <td></td>
1519 <td class="paramtype">uint&#160;</td>
1520 <td class="paramname"><em>max_step_x</em>, </td>
1521 </tr>
1522 <tr>
1523 <td class="paramkey"></td>
1524 <td></td>
1525 <td class="paramtype">uint&#160;</td>
1526 <td class="paramname"><em>max_stride_y</em>, </td>
1527 </tr>
1528 <tr>
1529 <td class="paramkey"></td>
1530 <td></td>
1531 <td class="paramtype">uint&#160;</td>
1532 <td class="paramname"><em>max_step_y</em>, </td>
1533 </tr>
1534 <tr>
1535 <td class="paramkey"></td>
1536 <td></td>
1537 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1538 <td class="paramname"><em>max_stride_z</em>, </td>
1539 </tr>
1540 <tr>
1541 <td class="paramkey"></td>
1542 <td></td>
1543 <td class="paramtype">uint&#160;</td>
1544 <td class="paramname"><em>max_step_z</em>, </td>
1545 </tr>
1546 <tr>
1547 <td class="paramkey"></td>
1548 <td></td>
1549 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1550 <td class="paramname"><em>max_offset_first_element_in_bytes</em>, </td>
1551 </tr>
1552 <tr>
1553 <td class="paramkey"></td>
1554 <td></td>
1555 <td class="paramtype">__global uchar *&#160;</td>
1556 <td class="paramname"><em>dst_ptr</em>, </td>
1557 </tr>
1558 <tr>
1559 <td class="paramkey"></td>
1560 <td></td>
1561 <td class="paramtype">uint&#160;</td>
1562 <td class="paramname"><em>dst_stride_x</em>, </td>
1563 </tr>
1564 <tr>
1565 <td class="paramkey"></td>
1566 <td></td>
1567 <td class="paramtype">uint&#160;</td>
1568 <td class="paramname"><em>dst_step_x</em>, </td>
1569 </tr>
1570 <tr>
1571 <td class="paramkey"></td>
1572 <td></td>
1573 <td class="paramtype">uint&#160;</td>
1574 <td class="paramname"><em>dst_stride_y</em>, </td>
1575 </tr>
1576 <tr>
1577 <td class="paramkey"></td>
1578 <td></td>
1579 <td class="paramtype">uint&#160;</td>
1580 <td class="paramname"><em>dst_step_y</em>, </td>
1581 </tr>
1582 <tr>
1583 <td class="paramkey"></td>
1584 <td></td>
1585 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1586 <td class="paramname"><em>dst_stride_z</em>, </td>
1587 </tr>
1588 <tr>
1589 <td class="paramkey"></td>
1590 <td></td>
1591 <td class="paramtype">uint&#160;</td>
1592 <td class="paramname"><em>dst_step_z</em>, </td>
1593 </tr>
1594 <tr>
1595 <td class="paramkey"></td>
1596 <td></td>
1597 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1598 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
1599 </tr>
1600 <tr>
1601 <td class="paramkey"></td>
1602 <td></td>
1603 <td class="paramtype">__global uchar *&#160;</td>
1604 <td class="paramname"><em>sum_ptr</em>, </td>
1605 </tr>
1606 <tr>
1607 <td class="paramkey"></td>
1608 <td></td>
1609 <td class="paramtype">uint&#160;</td>
1610 <td class="paramname"><em>sum_stride_x</em>, </td>
1611 </tr>
1612 <tr>
1613 <td class="paramkey"></td>
1614 <td></td>
1615 <td class="paramtype">uint&#160;</td>
1616 <td class="paramname"><em>sum_step_x</em>, </td>
1617 </tr>
1618 <tr>
1619 <td class="paramkey"></td>
1620 <td></td>
1621 <td class="paramtype">uint&#160;</td>
1622 <td class="paramname"><em>sum_stride_y</em>, </td>
1623 </tr>
1624 <tr>
1625 <td class="paramkey"></td>
1626 <td></td>
1627 <td class="paramtype">uint&#160;</td>
1628 <td class="paramname"><em>sum_step_y</em>, </td>
1629 </tr>
1630 <tr>
1631 <td class="paramkey"></td>
1632 <td></td>
1633 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1634 <td class="paramname"><em>sum_stride_z</em>, </td>
1635 </tr>
1636 <tr>
1637 <td class="paramkey"></td>
1638 <td></td>
1639 <td class="paramtype">uint&#160;</td>
1640 <td class="paramname"><em>sum_step_z</em>, </td>
1641 </tr>
1642 <tr>
1643 <td class="paramkey"></td>
1644 <td></td>
1645 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1646 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1647 </tr>
1648 <tr>
1649 <td class="paramkey"></td>
1650 <td></td>
1651 <td class="paramtype">uint&#160;</td>
1652 <td class="paramname"><em>width</em>&#160;</td>
1653 </tr>
1654 <tr>
1655 <td></td>
1656 <td>)</td>
1657 <td></td><td></td>
1658 </tr>
1659 </table>
1660</div><div class="memdoc">
1661
1662<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>
1663<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>
1664<dd>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1665Fixed 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]1666<dd>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1667In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed. </dd>
1668<dd>
1669Beta can be optionally passed at compile time using -DBETA (if undefined, assume it equals 1.0)</dd></dl>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1670<dl class="params"><dt>Parameters</dt><dd>
1671 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1672 <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]1673 <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>
1674 <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>
1675 <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>
1676 <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]1677 <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>
1678 <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]1679 <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]1680 <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]1681 <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>
1682 <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>
1683 <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>
1684 <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]1685 <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>
1686 <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]1687 <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]1688 <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]1689 <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>
1690 <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>
1691 <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>
1692 <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]1693 <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>
1694 <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]1695 <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]1696 <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]1697 <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>
1698 <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>
1699 <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]1700 <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>
1701 <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>
1702 <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]1703 <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>
1704 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1705 </table>
1706 </dd>
1707</dl>
1708
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1709<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00196">196</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]1710
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1711<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00112">arm_compute::test::validation::beta</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00123">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00048">EXP_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00087">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="softmax__layer_8cl_source.xhtml#l00046">MUL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">SELECT_DATA_TYPE</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00133">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">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>
1712<div class="fragment"><div class="line"><a name="l00202"></a><span class="lineno"> 202</span>&#160;{</div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00205"></a><span class="lineno"> 205</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><div class="line"><a name="l00206"></a><span class="lineno"> 206</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><div class="line"><a name="l00207"></a><span class="lineno"> 207</span>&#160;</div><div class="line"><a name="l00208"></a><span class="lineno"> 208</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00210"></a><span class="lineno"> 210</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><div class="line"><a name="l00211"></a><span class="lineno"> 211</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a> = (<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))BETA;</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>&#160;</div><div class="line"><a name="l00214"></a><span class="lineno"> 214</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>&#160; DATA_TYPE max_val = *((__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;max, 0, 0));</div><div class="line"><a name="l00216"></a><span class="lineno"> 216</span>&#160;</div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>&#160; <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span>&#160; sum1D = 0;</div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span>&#160;</div><div class="line"><a name="l00221"></a><span class="lineno"> 221</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00222"></a><span class="lineno"> 222</span>&#160; const uint width4 = width &gt;&gt; 4;</div><div class="line"><a name="l00223"></a><span class="lineno"> 223</span>&#160; for(uint i = 0; i &lt; width4; i++)</div><div class="line"><a name="l00224"></a><span class="lineno"> 224</span>&#160; {</div><div class="line"><a name="l00225"></a><span class="lineno"> 225</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00226"></a><span class="lineno"> 226</span>&#160; data = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, i &lt;&lt; 4, 0));</div><div class="line"><a name="l00227"></a><span class="lineno"> 227</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, 16);</div><div class="line"><a name="l00228"></a><span class="lineno"> 228</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00229"></a><span class="lineno"> 229</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a>, DATA_TYPE, 16);</div><div class="line"><a name="l00230"></a><span class="lineno"> 230</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00231"></a><span class="lineno"> 231</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, 16);</div><div class="line"><a name="l00232"></a><span class="lineno"> 232</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><div class="line"><a name="l00233"></a><span class="lineno"> 233</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, 16);</div><div class="line"><a name="l00234"></a><span class="lineno"> 234</span>&#160; }</div><div class="line"><a name="l00235"></a><span class="lineno"> 235</span>&#160;</div><div class="line"><a name="l00236"></a><span class="lineno"> 236</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_16</span></div><div class="line"><a name="l00237"></a><span class="lineno"> 237</span>&#160; <span class="comment">// Handle non multiple of 16</span></div><div class="line"><a name="l00238"></a><span class="lineno"> 238</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 16)</div><div class="line"><a name="l00239"></a><span class="lineno"> 239</span>&#160; data = vload16(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, width4 &lt;&lt; 4, 0));</div><div class="line"><a name="l00240"></a><span class="lineno"> 240</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, DATA_TYPE, 16);</div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00242"></a><span class="lineno"> 242</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a>, DATA_TYPE, 16);</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, DATA_TYPE, 16);</div><div class="line"><a name="l00245"></a><span class="lineno"> 245</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><div class="line"><a name="l00246"></a><span class="lineno"> 246</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>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 16));</div><div class="line"><a name="l00247"></a><span class="lineno"> 247</span>&#160; data = select(0, data, widx);</div><div class="line"><a name="l00248"></a><span class="lineno"> 248</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><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, DATA_TYPE, 16);</div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_16 */</span><span class="preprocessor"></span></div><div class="line"><a name="l00251"></a><span class="lineno"> 251</span>&#160;</div><div class="line"><a name="l00252"></a><span class="lineno"> 252</span>&#160; <span class="comment">// Perform min/max reduction</span></div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span>&#160; sum1D.s01234567 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01234567, sum1D.s89ABCDEF, DATA_TYPE, 8);</div><div class="line"><a name="l00254"></a><span class="lineno"> 254</span>&#160; sum1D.s0123 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0123, sum1D.s4567, DATA_TYPE, 4);</div><div class="line"><a name="l00255"></a><span class="lineno"> 255</span>&#160; sum1D.s01 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01, sum1D.s23, DATA_TYPE, 2);</div><div class="line"><a name="l00256"></a><span class="lineno"> 256</span>&#160; sum1D.s0 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0, sum1D.s1, DATA_TYPE, 1);</div><div class="line"><a name="l00257"></a><span class="lineno"> 257</span>&#160;</div><div class="line"><a name="l00258"></a><span class="lineno"> 258</span>&#160; <span class="comment">// Calculate and store result</span></div><div class="line"><a name="l00259"></a><span class="lineno"> 259</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div><div class="line"><a name="l00260"></a><span class="lineno"> 260</span>&#160;}</div><div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ab146b9cbab6e73e7588b240dc709fe01"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">arm_compute::test::validation::beta</a></div><div class="ttdeci">beta</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00112">GEMM.cpp:112</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1713<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1714<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#l00048">softmax_layer.cl:48</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1715<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>
1716<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1717<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#l00044">softmax_layer.cl:44</a></div></div>
1718<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
1719<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#l00055">softmax_layer.cl:55</a></div></div>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1720<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00123">Convolution.cpp:123</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1721<div class="ttc" id="softmax__layer_8cl_xhtml_a22303c4047ec5027c1538d53964b9d0d"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00046">softmax_layer.cl:46</a></div></div>
1722<div class="ttc" id="struct_image_xhtml"><div class="ttname"><a href="struct_image.xhtml">Image</a></div><div class="ttdoc">Structure to hold Image information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
1723<div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00142">helpers.h:142</a></div></div>
1724<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#l00045">softmax_layer.cl:45</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1725<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1726<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#l00087">softmax_layer.cl:87</a></div></div>
1727<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1728<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 Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1729<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00133">Convolution.cpp:133</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1730</div><!-- fragment -->
1731</div>
1732</div>
1733<h2 class="groupheader">Variable Documentation</h2>
1734<a class="anchor" id="a0712735973f172ac9efc7d48a31e47ad"></a>
1735<div class="memitem">
1736<div class="memproto">
1737 <table class="memname">
1738 <tr>
1739 <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1740 </tr>
1741 </table>
1742</div><div class="memdoc">
1743
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1744<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00087">87</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]1745
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1746<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00114">softmax_layer_max()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00196">softmax_layer_shift_exp_sum()</a>.</p>
1747
1748</div>
1749</div>
1750<a class="anchor" id="a4884a666a1e93fbf8c27bd7d2da3c8bb"></a>
1751<div class="memitem">
1752<div class="memproto">
1753 <table class="memname">
1754 <tr>
1755 <td class="memname">__constant uint4 idx4 = (uint4)(0, 1, 2, 3)</td>
1756 </tr>
1757 </table>
1758</div><div class="memdoc">
1759
1760<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00088">88</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1761
1762<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>.</p>
1763
1764</div>
1765</div>
1766<a class="anchor" id="aa1dd94b8d98f1c6d790bdf0fc5de29e9"></a>
1767<div class="memitem">
1768<div class="memproto">
1769 <table class="memname">
1770 <tr>
1771 <td class="memname">__constant uint16 idx__ = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1772 </tr>
1773 </table>
1774</div><div class="memdoc">
1775
1776<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00082">82</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1777
1778<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1779
1780</div>
1781</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1782<a class="anchor" id="a538b4b63f40e7b12891774e03a4f0dec"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1783<div class="memitem">
1784<div class="memproto">
1785 <table class="memname">
1786 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1787 <td class="memname">__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( -FLT_MAX )</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1788 </tr>
1789 </table>
1790</div><div class="memdoc">
1791
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1792<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00086">86</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]1793
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1794<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00114">softmax_layer_max()</a>.</p>
1795
1796</div>
1797</div>
1798<a class="anchor" id="af7a08044d0e491a0ee1520a24a107a2b"></a>
1799<div class="memitem">
1800<div class="memproto">
1801 <table class="memname">
1802 <tr>
1803 <td class="memname">__constant DATA_TYPE16 type_min_ = ( DATA_TYPE16 )( -FLT_MAX )</td>
1804 </tr>
1805 </table>
1806</div><div class="memdoc">
1807
1808<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00081">81</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1809
1810<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00350">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1811
1812</div>
1813</div>
1814</div><!-- contents -->
1815</div><!-- doc-content -->
1816<!-- start footer part -->
1817<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
1818 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1819 <li class="navelem"><a class="el" href="dir_68267d1309a1af8e8297ef4c3efbcdba.xhtml">src</a></li><li class="navelem"><a class="el" href="dir_aebb8dcc11953d78e620bbef0b9e2183.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_8c278f79c760e5c5fbd911f9870614c1.xhtml">CL</a></li><li class="navelem"><a class="el" href="dir_25885286e9dad4fa105b7b25a8031bbf.xhtml">cl_kernels</a></li><li class="navelem"><a class="el" href="softmax__layer_8cl.xhtml">softmax_layer.cl</a></li>
Anthony Barbierf45d5a92018-01-24 16:23:15 +0000[diff] [blame]1820 <li class="footer">Generated on Wed Jan 24 2018 14:30:43 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1821 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1822 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1823 </ul>
1824</div>
1825</body>
1826</html>