Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 8140e1e155d3430992fa46e04ef8938ff09ffd2d / . / documentation / softmax__layer_8cl.xhtml
blob: e989f2274cdbfca599ac252cf8ef6502756f3c92 [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 Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]43 &#160;<span id="projectnumber">17.12</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>
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> src = <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>
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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
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>
668</div><!-- fragment -->
669</div>
670</div>
671<a class="anchor" id="a82bc360279fcf7cf8033dca6022f21f5"></a>
672<div class="memitem">
673<div class="memproto">
674 <table class="memname">
675 <tr>
676 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_parallel </td>
677 <td>(</td>
678 <td class="paramtype">__global uchar *&#160;</td>
679 <td class="paramname"><em>src_ptr</em>, </td>
680 </tr>
681 <tr>
682 <td class="paramkey"></td>
683 <td></td>
684 <td class="paramtype">uint&#160;</td>
685 <td class="paramname"><em>src_stride_x</em>, </td>
686 </tr>
687 <tr>
688 <td class="paramkey"></td>
689 <td></td>
690 <td class="paramtype">uint&#160;</td>
691 <td class="paramname"><em>src_step_x</em>, </td>
692 </tr>
693 <tr>
694 <td class="paramkey"></td>
695 <td></td>
696 <td class="paramtype">uint&#160;</td>
697 <td class="paramname"><em>src_stride_y</em>, </td>
698 </tr>
699 <tr>
700 <td class="paramkey"></td>
701 <td></td>
702 <td class="paramtype">uint&#160;</td>
703 <td class="paramname"><em>src_step_y</em>, </td>
704 </tr>
705 <tr>
706 <td class="paramkey"></td>
707 <td></td>
708 <td class="paramtype">uint&#160;</td>
709 <td class="paramname"><em>src_stride_z</em>, </td>
710 </tr>
711 <tr>
712 <td class="paramkey"></td>
713 <td></td>
714 <td class="paramtype">uint&#160;</td>
715 <td class="paramname"><em>src_step_z</em>, </td>
716 </tr>
717 <tr>
718 <td class="paramkey"></td>
719 <td></td>
720 <td class="paramtype">uint&#160;</td>
721 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
722 </tr>
723 <tr>
724 <td class="paramkey"></td>
725 <td></td>
726 <td class="paramtype">__global uchar *&#160;</td>
727 <td class="paramname"><em>maxo_ptr</em>, </td>
728 </tr>
729 <tr>
730 <td class="paramkey"></td>
731 <td></td>
732 <td class="paramtype">uint&#160;</td>
733 <td class="paramname"><em>maxo_stride_x</em>, </td>
734 </tr>
735 <tr>
736 <td class="paramkey"></td>
737 <td></td>
738 <td class="paramtype">uint&#160;</td>
739 <td class="paramname"><em>maxo_step_x</em>, </td>
740 </tr>
741 <tr>
742 <td class="paramkey"></td>
743 <td></td>
744 <td class="paramtype">uint&#160;</td>
745 <td class="paramname"><em>maxo_stride_y</em>, </td>
746 </tr>
747 <tr>
748 <td class="paramkey"></td>
749 <td></td>
750 <td class="paramtype">uint&#160;</td>
751 <td class="paramname"><em>maxo_step_y</em>, </td>
752 </tr>
753 <tr>
754 <td class="paramkey"></td>
755 <td></td>
756 <td class="paramtype">uint&#160;</td>
757 <td class="paramname"><em>maxo_stride_z</em>, </td>
758 </tr>
759 <tr>
760 <td class="paramkey"></td>
761 <td></td>
762 <td class="paramtype">uint&#160;</td>
763 <td class="paramname"><em>maxo_step_z</em>, </td>
764 </tr>
765 <tr>
766 <td class="paramkey"></td>
767 <td></td>
768 <td class="paramtype">uint&#160;</td>
769 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
770 </tr>
771 <tr>
772 <td class="paramkey"></td>
773 <td></td>
774 <td class="paramtype">__global uchar *&#160;</td>
775 <td class="paramname"><em>dst_ptr</em>, </td>
776 </tr>
777 <tr>
778 <td class="paramkey"></td>
779 <td></td>
780 <td class="paramtype">uint&#160;</td>
781 <td class="paramname"><em>dst_stride_x</em>, </td>
782 </tr>
783 <tr>
784 <td class="paramkey"></td>
785 <td></td>
786 <td class="paramtype">uint&#160;</td>
787 <td class="paramname"><em>dst_step_x</em>, </td>
788 </tr>
789 <tr>
790 <td class="paramkey"></td>
791 <td></td>
792 <td class="paramtype">uint&#160;</td>
793 <td class="paramname"><em>dst_stride_y</em>, </td>
794 </tr>
795 <tr>
796 <td class="paramkey"></td>
797 <td></td>
798 <td class="paramtype">uint&#160;</td>
799 <td class="paramname"><em>dst_step_y</em>, </td>
800 </tr>
801 <tr>
802 <td class="paramkey"></td>
803 <td></td>
804 <td class="paramtype">uint&#160;</td>
805 <td class="paramname"><em>dst_stride_z</em>, </td>
806 </tr>
807 <tr>
808 <td class="paramkey"></td>
809 <td></td>
810 <td class="paramtype">uint&#160;</td>
811 <td class="paramname"><em>dst_step_z</em>, </td>
812 </tr>
813 <tr>
814 <td class="paramkey"></td>
815 <td></td>
816 <td class="paramtype">uint&#160;</td>
817 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
818 </tr>
819 <tr>
820 <td class="paramkey"></td>
821 <td></td>
822 <td class="paramtype">__global uchar *&#160;</td>
823 <td class="paramname"><em>sum_ptr</em>, </td>
824 </tr>
825 <tr>
826 <td class="paramkey"></td>
827 <td></td>
828 <td class="paramtype">uint&#160;</td>
829 <td class="paramname"><em>sum_stride_x</em>, </td>
830 </tr>
831 <tr>
832 <td class="paramkey"></td>
833 <td></td>
834 <td class="paramtype">uint&#160;</td>
835 <td class="paramname"><em>sum_step_x</em>, </td>
836 </tr>
837 <tr>
838 <td class="paramkey"></td>
839 <td></td>
840 <td class="paramtype">uint&#160;</td>
841 <td class="paramname"><em>sum_stride_y</em>, </td>
842 </tr>
843 <tr>
844 <td class="paramkey"></td>
845 <td></td>
846 <td class="paramtype">uint&#160;</td>
847 <td class="paramname"><em>sum_step_y</em>, </td>
848 </tr>
849 <tr>
850 <td class="paramkey"></td>
851 <td></td>
852 <td class="paramtype">uint&#160;</td>
853 <td class="paramname"><em>sum_stride_z</em>, </td>
854 </tr>
855 <tr>
856 <td class="paramkey"></td>
857 <td></td>
858 <td class="paramtype">uint&#160;</td>
859 <td class="paramname"><em>sum_step_z</em>, </td>
860 </tr>
861 <tr>
862 <td class="paramkey"></td>
863 <td></td>
864 <td class="paramtype">uint&#160;</td>
865 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
866 </tr>
867 <tr>
868 <td class="paramkey"></td>
869 <td></td>
870 <td class="paramtype">uint&#160;</td>
871 <td class="paramname"><em>width</em>&#160;</td>
872 </tr>
873 <tr>
874 <td></td>
875 <td>)</td>
876 <td></td><td></td>
877 </tr>
878 </table>
879</div><div class="memdoc">
880
881<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>
882<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>
883<dd>
884Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
885<dd>
886In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
887<dd>
888Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
889<dl class="params"><dt>Parameters</dt><dd>
890 <table class="params">
891 <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>
892 <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>
893 <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>
894 <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>
895 <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>
896 <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>
897 <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>
898 <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>
899 <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>
900 <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>
901 <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>
902 <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>
903 <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>
904 <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>
905 <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>
906 <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>
907 <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>
908 <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>
909 <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>
910 <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>
911 <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>
912 <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>
913 <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>
914 <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>
915 <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>
916 <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>
917 <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>
918 <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>
919 <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>
920 <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>
921 <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>
922 <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>
923 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
924 </table>
925 </dd>
926</dl>
927
928<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>
929
930<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00109">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="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">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="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>
931<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> src = <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_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00109">GEMM.cpp:109</a></div></div>
932<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>
933<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>
934<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>
935<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>
936<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>
937<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>
938<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>
939<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>
940<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>
941<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>
942<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>
943<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>
944<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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
945<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>
946<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>
947<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>
948<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>
949<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>
950<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>
951<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>
952<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>
953</div><!-- fragment -->
954</div>
955</div>
956<a class="anchor" id="a9d74601bcabbb2f14bcf52385ad666dc"></a>
957<div class="memitem">
958<div class="memproto">
959 <table class="memname">
960 <tr>
961 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_serial </td>
962 <td>(</td>
963 <td class="paramtype">__global uchar *&#160;</td>
964 <td class="paramname"><em>src_ptr</em>, </td>
965 </tr>
966 <tr>
967 <td class="paramkey"></td>
968 <td></td>
969 <td class="paramtype">uint&#160;</td>
970 <td class="paramname"><em>src_stride_x</em>, </td>
971 </tr>
972 <tr>
973 <td class="paramkey"></td>
974 <td></td>
975 <td class="paramtype">uint&#160;</td>
976 <td class="paramname"><em>src_step_x</em>, </td>
977 </tr>
978 <tr>
979 <td class="paramkey"></td>
980 <td></td>
981 <td class="paramtype">uint&#160;</td>
982 <td class="paramname"><em>src_stride_y</em>, </td>
983 </tr>
984 <tr>
985 <td class="paramkey"></td>
986 <td></td>
987 <td class="paramtype">uint&#160;</td>
988 <td class="paramname"><em>src_step_y</em>, </td>
989 </tr>
990 <tr>
991 <td class="paramkey"></td>
992 <td></td>
993 <td class="paramtype">uint&#160;</td>
994 <td class="paramname"><em>src_stride_z</em>, </td>
995 </tr>
996 <tr>
997 <td class="paramkey"></td>
998 <td></td>
999 <td class="paramtype">uint&#160;</td>
1000 <td class="paramname"><em>src_step_z</em>, </td>
1001 </tr>
1002 <tr>
1003 <td class="paramkey"></td>
1004 <td></td>
1005 <td class="paramtype">uint&#160;</td>
1006 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1007 </tr>
1008 <tr>
1009 <td class="paramkey"></td>
1010 <td></td>
1011 <td class="paramtype">__global uchar *&#160;</td>
1012 <td class="paramname"><em>maxo_ptr</em>, </td>
1013 </tr>
1014 <tr>
1015 <td class="paramkey"></td>
1016 <td></td>
1017 <td class="paramtype">uint&#160;</td>
1018 <td class="paramname"><em>maxo_stride_x</em>, </td>
1019 </tr>
1020 <tr>
1021 <td class="paramkey"></td>
1022 <td></td>
1023 <td class="paramtype">uint&#160;</td>
1024 <td class="paramname"><em>maxo_step_x</em>, </td>
1025 </tr>
1026 <tr>
1027 <td class="paramkey"></td>
1028 <td></td>
1029 <td class="paramtype">uint&#160;</td>
1030 <td class="paramname"><em>maxo_stride_y</em>, </td>
1031 </tr>
1032 <tr>
1033 <td class="paramkey"></td>
1034 <td></td>
1035 <td class="paramtype">uint&#160;</td>
1036 <td class="paramname"><em>maxo_step_y</em>, </td>
1037 </tr>
1038 <tr>
1039 <td class="paramkey"></td>
1040 <td></td>
1041 <td class="paramtype">uint&#160;</td>
1042 <td class="paramname"><em>maxo_stride_z</em>, </td>
1043 </tr>
1044 <tr>
1045 <td class="paramkey"></td>
1046 <td></td>
1047 <td class="paramtype">uint&#160;</td>
1048 <td class="paramname"><em>maxo_step_z</em>, </td>
1049 </tr>
1050 <tr>
1051 <td class="paramkey"></td>
1052 <td></td>
1053 <td class="paramtype">uint&#160;</td>
1054 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
1055 </tr>
1056 <tr>
1057 <td class="paramkey"></td>
1058 <td></td>
1059 <td class="paramtype">__global uchar *&#160;</td>
1060 <td class="paramname"><em>dst_ptr</em>, </td>
1061 </tr>
1062 <tr>
1063 <td class="paramkey"></td>
1064 <td></td>
1065 <td class="paramtype">uint&#160;</td>
1066 <td class="paramname"><em>dst_stride_x</em>, </td>
1067 </tr>
1068 <tr>
1069 <td class="paramkey"></td>
1070 <td></td>
1071 <td class="paramtype">uint&#160;</td>
1072 <td class="paramname"><em>dst_step_x</em>, </td>
1073 </tr>
1074 <tr>
1075 <td class="paramkey"></td>
1076 <td></td>
1077 <td class="paramtype">uint&#160;</td>
1078 <td class="paramname"><em>dst_stride_y</em>, </td>
1079 </tr>
1080 <tr>
1081 <td class="paramkey"></td>
1082 <td></td>
1083 <td class="paramtype">uint&#160;</td>
1084 <td class="paramname"><em>dst_step_y</em>, </td>
1085 </tr>
1086 <tr>
1087 <td class="paramkey"></td>
1088 <td></td>
1089 <td class="paramtype">uint&#160;</td>
1090 <td class="paramname"><em>dst_stride_z</em>, </td>
1091 </tr>
1092 <tr>
1093 <td class="paramkey"></td>
1094 <td></td>
1095 <td class="paramtype">uint&#160;</td>
1096 <td class="paramname"><em>dst_step_z</em>, </td>
1097 </tr>
1098 <tr>
1099 <td class="paramkey"></td>
1100 <td></td>
1101 <td class="paramtype">uint&#160;</td>
1102 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
1103 </tr>
1104 <tr>
1105 <td class="paramkey"></td>
1106 <td></td>
1107 <td class="paramtype">__global uchar *&#160;</td>
1108 <td class="paramname"><em>sum_ptr</em>, </td>
1109 </tr>
1110 <tr>
1111 <td class="paramkey"></td>
1112 <td></td>
1113 <td class="paramtype">uint&#160;</td>
1114 <td class="paramname"><em>sum_stride_x</em>, </td>
1115 </tr>
1116 <tr>
1117 <td class="paramkey"></td>
1118 <td></td>
1119 <td class="paramtype">uint&#160;</td>
1120 <td class="paramname"><em>sum_step_x</em>, </td>
1121 </tr>
1122 <tr>
1123 <td class="paramkey"></td>
1124 <td></td>
1125 <td class="paramtype">uint&#160;</td>
1126 <td class="paramname"><em>sum_stride_y</em>, </td>
1127 </tr>
1128 <tr>
1129 <td class="paramkey"></td>
1130 <td></td>
1131 <td class="paramtype">uint&#160;</td>
1132 <td class="paramname"><em>sum_step_y</em>, </td>
1133 </tr>
1134 <tr>
1135 <td class="paramkey"></td>
1136 <td></td>
1137 <td class="paramtype">uint&#160;</td>
1138 <td class="paramname"><em>sum_stride_z</em>, </td>
1139 </tr>
1140 <tr>
1141 <td class="paramkey"></td>
1142 <td></td>
1143 <td class="paramtype">uint&#160;</td>
1144 <td class="paramname"><em>sum_step_z</em>, </td>
1145 </tr>
1146 <tr>
1147 <td class="paramkey"></td>
1148 <td></td>
1149 <td class="paramtype">uint&#160;</td>
1150 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1151 </tr>
1152 <tr>
1153 <td class="paramkey"></td>
1154 <td></td>
1155 <td class="paramtype">uint&#160;</td>
1156 <td class="paramname"><em>width</em>&#160;</td>
1157 </tr>
1158 <tr>
1159 <td></td>
1160 <td>)</td>
1161 <td></td><td></td>
1162 </tr>
1163 </table>
1164</div><div class="memdoc">
1165
1166<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>
1167<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>
1168<dd>
1169Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
1170<dd>
1171In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
1172<dd>
1173Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
1174<dl class="params"><dt>Parameters</dt><dd>
1175 <table class="params">
1176 <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>
1177 <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>
1178 <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>
1179 <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>
1180 <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>
1181 <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>
1182 <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>
1183 <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>
1184 <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>
1185 <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>
1186 <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>
1187 <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>
1188 <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>
1189 <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>
1190 <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>
1191 <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>
1192 <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>
1193 <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>
1194 <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>
1195 <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>
1196 <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>
1197 <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>
1198 <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>
1199 <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>
1200 <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>
1201 <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>
1202 <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>
1203 <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>
1204 <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>
1205 <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>
1206 <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>
1207 <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>
1208 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1209 </table>
1210 </dd>
1211</dl>
1212
1213<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>
1214
1215<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00109">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="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">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="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>
1216<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> src = <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>
1217<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_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00109">GEMM.cpp:109</a></div></div>
1218<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>
1219<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>
1220<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>
1221<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>
1222<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>
1223<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>
1224<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>
1225<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>
1226<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>
1227<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>
1228<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>
1229<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>
1230<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>
1231<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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
1232<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>
1233<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>
1234<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>
1235<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>
1236<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>
1237<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>
1238<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>
1239<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 Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1240</div><!-- fragment -->
1241</div>
1242</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1243<a class="anchor" id="ac4247ac0991e85965b7ded764e78f12c"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1244<div class="memitem">
1245<div class="memproto">
1246 <table class="memname">
1247 <tr>
1248 <td class="memname">__kernel void softmax_layer_norm </td>
1249 <td>(</td>
1250 <td class="paramtype">__global uchar *&#160;</td>
1251 <td class="paramname"><em>src_ptr</em>, </td>
1252 </tr>
1253 <tr>
1254 <td class="paramkey"></td>
1255 <td></td>
1256 <td class="paramtype">uint&#160;</td>
1257 <td class="paramname"><em>src_stride_x</em>, </td>
1258 </tr>
1259 <tr>
1260 <td class="paramkey"></td>
1261 <td></td>
1262 <td class="paramtype">uint&#160;</td>
1263 <td class="paramname"><em>src_step_x</em>, </td>
1264 </tr>
1265 <tr>
1266 <td class="paramkey"></td>
1267 <td></td>
1268 <td class="paramtype">uint&#160;</td>
1269 <td class="paramname"><em>src_stride_y</em>, </td>
1270 </tr>
1271 <tr>
1272 <td class="paramkey"></td>
1273 <td></td>
1274 <td class="paramtype">uint&#160;</td>
1275 <td class="paramname"><em>src_step_y</em>, </td>
1276 </tr>
1277 <tr>
1278 <td class="paramkey"></td>
1279 <td></td>
1280 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1281 <td class="paramname"><em>src_stride_z</em>, </td>
1282 </tr>
1283 <tr>
1284 <td class="paramkey"></td>
1285 <td></td>
1286 <td class="paramtype">uint&#160;</td>
1287 <td class="paramname"><em>src_step_z</em>, </td>
1288 </tr>
1289 <tr>
1290 <td class="paramkey"></td>
1291 <td></td>
1292 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1293 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1294 </tr>
1295 <tr>
1296 <td class="paramkey"></td>
1297 <td></td>
1298 <td class="paramtype">__global uchar *&#160;</td>
1299 <td class="paramname"><em>sum_ptr</em>, </td>
1300 </tr>
1301 <tr>
1302 <td class="paramkey"></td>
1303 <td></td>
1304 <td class="paramtype">uint&#160;</td>
1305 <td class="paramname"><em>sum_stride_x</em>, </td>
1306 </tr>
1307 <tr>
1308 <td class="paramkey"></td>
1309 <td></td>
1310 <td class="paramtype">uint&#160;</td>
1311 <td class="paramname"><em>sum_step_x</em>, </td>
1312 </tr>
1313 <tr>
1314 <td class="paramkey"></td>
1315 <td></td>
1316 <td class="paramtype">uint&#160;</td>
1317 <td class="paramname"><em>sum_stride_y</em>, </td>
1318 </tr>
1319 <tr>
1320 <td class="paramkey"></td>
1321 <td></td>
1322 <td class="paramtype">uint&#160;</td>
1323 <td class="paramname"><em>sum_step_y</em>, </td>
1324 </tr>
1325 <tr>
1326 <td class="paramkey"></td>
1327 <td></td>
1328 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1329 <td class="paramname"><em>sum_stride_z</em>, </td>
1330 </tr>
1331 <tr>
1332 <td class="paramkey"></td>
1333 <td></td>
1334 <td class="paramtype">uint&#160;</td>
1335 <td class="paramname"><em>sum_step_z</em>, </td>
1336 </tr>
1337 <tr>
1338 <td class="paramkey"></td>
1339 <td></td>
1340 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1341 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1342 </tr>
1343 <tr>
1344 <td class="paramkey"></td>
1345 <td></td>
1346 <td class="paramtype">__global uchar *&#160;</td>
1347 <td class="paramname"><em>dst_ptr</em>, </td>
1348 </tr>
1349 <tr>
1350 <td class="paramkey"></td>
1351 <td></td>
1352 <td class="paramtype">uint&#160;</td>
1353 <td class="paramname"><em>dst_stride_x</em>, </td>
1354 </tr>
1355 <tr>
1356 <td class="paramkey"></td>
1357 <td></td>
1358 <td class="paramtype">uint&#160;</td>
1359 <td class="paramname"><em>dst_step_x</em>, </td>
1360 </tr>
1361 <tr>
1362 <td class="paramkey"></td>
1363 <td></td>
1364 <td class="paramtype">uint&#160;</td>
1365 <td class="paramname"><em>dst_stride_y</em>, </td>
1366 </tr>
1367 <tr>
1368 <td class="paramkey"></td>
1369 <td></td>
1370 <td class="paramtype">uint&#160;</td>
1371 <td class="paramname"><em>dst_step_y</em>, </td>
1372 </tr>
1373 <tr>
1374 <td class="paramkey"></td>
1375 <td></td>
1376 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1377 <td class="paramname"><em>dst_stride_z</em>, </td>
1378 </tr>
1379 <tr>
1380 <td class="paramkey"></td>
1381 <td></td>
1382 <td class="paramtype">uint&#160;</td>
1383 <td class="paramname"><em>dst_step_z</em>, </td>
1384 </tr>
1385 <tr>
1386 <td class="paramkey"></td>
1387 <td></td>
1388 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1389 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
1390 </tr>
1391 <tr>
1392 <td></td>
1393 <td>)</td>
1394 <td></td><td></td>
1395 </tr>
1396 </table>
1397</div><div class="memdoc">
1398
1399<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]1400<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>
1401<dd>
1402Fixed 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]1403<dl class="params"><dt>Parameters</dt><dd>
1404 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1405 <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]1406 <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>
1407 <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>
1408 <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>
1409 <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]1410 <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>
1411 <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]1412 <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]1413 <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]1414 <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>
1415 <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>
1416 <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>
1417 <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]1418 <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>
1419 <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]1420 <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]1421 <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]1422 <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>
1423 <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>
1424 <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>
1425 <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]1426 <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>
1427 <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]1428 <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>
1429 </table>
1430 </dd>
1431</dl>
1432
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1433<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]1434
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1435<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="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">arm_compute::test::validation::dst</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, and <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>.</p>
1436<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> src = <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]1437<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>
1438<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^]1439<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>
1440<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>
1441<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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
1442<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]1443<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^]1444<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 Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1445</div><!-- fragment -->
1446</div>
1447</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1448<a class="anchor" id="aa4412bb319c67deaeab07a481d93dfdb"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1449<div class="memitem">
1450<div class="memproto">
1451 <table class="memname">
1452 <tr>
1453 <td class="memname">__kernel void softmax_layer_shift_exp_sum </td>
1454 <td>(</td>
1455 <td class="paramtype">__global uchar *&#160;</td>
1456 <td class="paramname"><em>src_ptr</em>, </td>
1457 </tr>
1458 <tr>
1459 <td class="paramkey"></td>
1460 <td></td>
1461 <td class="paramtype">uint&#160;</td>
1462 <td class="paramname"><em>src_stride_x</em>, </td>
1463 </tr>
1464 <tr>
1465 <td class="paramkey"></td>
1466 <td></td>
1467 <td class="paramtype">uint&#160;</td>
1468 <td class="paramname"><em>src_step_x</em>, </td>
1469 </tr>
1470 <tr>
1471 <td class="paramkey"></td>
1472 <td></td>
1473 <td class="paramtype">uint&#160;</td>
1474 <td class="paramname"><em>src_stride_y</em>, </td>
1475 </tr>
1476 <tr>
1477 <td class="paramkey"></td>
1478 <td></td>
1479 <td class="paramtype">uint&#160;</td>
1480 <td class="paramname"><em>src_step_y</em>, </td>
1481 </tr>
1482 <tr>
1483 <td class="paramkey"></td>
1484 <td></td>
1485 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1486 <td class="paramname"><em>src_stride_z</em>, </td>
1487 </tr>
1488 <tr>
1489 <td class="paramkey"></td>
1490 <td></td>
1491 <td class="paramtype">uint&#160;</td>
1492 <td class="paramname"><em>src_step_z</em>, </td>
1493 </tr>
1494 <tr>
1495 <td class="paramkey"></td>
1496 <td></td>
1497 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1498 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1499 </tr>
1500 <tr>
1501 <td class="paramkey"></td>
1502 <td></td>
1503 <td class="paramtype">__global uchar *&#160;</td>
1504 <td class="paramname"><em>max_ptr</em>, </td>
1505 </tr>
1506 <tr>
1507 <td class="paramkey"></td>
1508 <td></td>
1509 <td class="paramtype">uint&#160;</td>
1510 <td class="paramname"><em>max_stride_x</em>, </td>
1511 </tr>
1512 <tr>
1513 <td class="paramkey"></td>
1514 <td></td>
1515 <td class="paramtype">uint&#160;</td>
1516 <td class="paramname"><em>max_step_x</em>, </td>
1517 </tr>
1518 <tr>
1519 <td class="paramkey"></td>
1520 <td></td>
1521 <td class="paramtype">uint&#160;</td>
1522 <td class="paramname"><em>max_stride_y</em>, </td>
1523 </tr>
1524 <tr>
1525 <td class="paramkey"></td>
1526 <td></td>
1527 <td class="paramtype">uint&#160;</td>
1528 <td class="paramname"><em>max_step_y</em>, </td>
1529 </tr>
1530 <tr>
1531 <td class="paramkey"></td>
1532 <td></td>
1533 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1534 <td class="paramname"><em>max_stride_z</em>, </td>
1535 </tr>
1536 <tr>
1537 <td class="paramkey"></td>
1538 <td></td>
1539 <td class="paramtype">uint&#160;</td>
1540 <td class="paramname"><em>max_step_z</em>, </td>
1541 </tr>
1542 <tr>
1543 <td class="paramkey"></td>
1544 <td></td>
1545 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1546 <td class="paramname"><em>max_offset_first_element_in_bytes</em>, </td>
1547 </tr>
1548 <tr>
1549 <td class="paramkey"></td>
1550 <td></td>
1551 <td class="paramtype">__global uchar *&#160;</td>
1552 <td class="paramname"><em>dst_ptr</em>, </td>
1553 </tr>
1554 <tr>
1555 <td class="paramkey"></td>
1556 <td></td>
1557 <td class="paramtype">uint&#160;</td>
1558 <td class="paramname"><em>dst_stride_x</em>, </td>
1559 </tr>
1560 <tr>
1561 <td class="paramkey"></td>
1562 <td></td>
1563 <td class="paramtype">uint&#160;</td>
1564 <td class="paramname"><em>dst_step_x</em>, </td>
1565 </tr>
1566 <tr>
1567 <td class="paramkey"></td>
1568 <td></td>
1569 <td class="paramtype">uint&#160;</td>
1570 <td class="paramname"><em>dst_stride_y</em>, </td>
1571 </tr>
1572 <tr>
1573 <td class="paramkey"></td>
1574 <td></td>
1575 <td class="paramtype">uint&#160;</td>
1576 <td class="paramname"><em>dst_step_y</em>, </td>
1577 </tr>
1578 <tr>
1579 <td class="paramkey"></td>
1580 <td></td>
1581 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1582 <td class="paramname"><em>dst_stride_z</em>, </td>
1583 </tr>
1584 <tr>
1585 <td class="paramkey"></td>
1586 <td></td>
1587 <td class="paramtype">uint&#160;</td>
1588 <td class="paramname"><em>dst_step_z</em>, </td>
1589 </tr>
1590 <tr>
1591 <td class="paramkey"></td>
1592 <td></td>
1593 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1594 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
1595 </tr>
1596 <tr>
1597 <td class="paramkey"></td>
1598 <td></td>
1599 <td class="paramtype">__global uchar *&#160;</td>
1600 <td class="paramname"><em>sum_ptr</em>, </td>
1601 </tr>
1602 <tr>
1603 <td class="paramkey"></td>
1604 <td></td>
1605 <td class="paramtype">uint&#160;</td>
1606 <td class="paramname"><em>sum_stride_x</em>, </td>
1607 </tr>
1608 <tr>
1609 <td class="paramkey"></td>
1610 <td></td>
1611 <td class="paramtype">uint&#160;</td>
1612 <td class="paramname"><em>sum_step_x</em>, </td>
1613 </tr>
1614 <tr>
1615 <td class="paramkey"></td>
1616 <td></td>
1617 <td class="paramtype">uint&#160;</td>
1618 <td class="paramname"><em>sum_stride_y</em>, </td>
1619 </tr>
1620 <tr>
1621 <td class="paramkey"></td>
1622 <td></td>
1623 <td class="paramtype">uint&#160;</td>
1624 <td class="paramname"><em>sum_step_y</em>, </td>
1625 </tr>
1626 <tr>
1627 <td class="paramkey"></td>
1628 <td></td>
1629 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1630 <td class="paramname"><em>sum_stride_z</em>, </td>
1631 </tr>
1632 <tr>
1633 <td class="paramkey"></td>
1634 <td></td>
1635 <td class="paramtype">uint&#160;</td>
1636 <td class="paramname"><em>sum_step_z</em>, </td>
1637 </tr>
1638 <tr>
1639 <td class="paramkey"></td>
1640 <td></td>
1641 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1642 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1643 </tr>
1644 <tr>
1645 <td class="paramkey"></td>
1646 <td></td>
1647 <td class="paramtype">uint&#160;</td>
1648 <td class="paramname"><em>width</em>&#160;</td>
1649 </tr>
1650 <tr>
1651 <td></td>
1652 <td>)</td>
1653 <td></td><td></td>
1654 </tr>
1655 </table>
1656</div><div class="memdoc">
1657
1658<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>
1659<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>
1660<dd>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1661Fixed 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]1662<dd>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1663In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed. </dd>
1664<dd>
1665Beta 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]1666<dl class="params"><dt>Parameters</dt><dd>
1667 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1668 <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]1669 <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>
1670 <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>
1671 <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>
1672 <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]1673 <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>
1674 <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]1675 <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]1676 <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]1677 <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>
1678 <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>
1679 <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>
1680 <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]1681 <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>
1682 <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]1683 <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]1684 <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]1685 <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>
1686 <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>
1687 <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>
1688 <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]1689 <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>
1690 <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]1691 <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]1692 <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]1693 <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>
1694 <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>
1695 <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]1696 <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>
1697 <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>
1698 <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]1699 <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>
1700 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1701 </table>
1702 </dd>
1703</dl>
1704
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1705<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]1706
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1707<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00109">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="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">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="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>
1708<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> src = <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_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00109">GEMM.cpp:109</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1709<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^]1710<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]1711<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>
1712<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^]1713<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>
1714<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>
1715<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>
1716<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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
1717<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>
1718<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>
1719<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>
1720<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]1721<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^]1722<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>
1723<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]1724<div class="ttc" id="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail_xhtml_ad91bb73431b4de1f4946ed949d444849"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1fixed__point__arithmetic_1_1detail.xhtml#ad91bb73431b4de1f4946ed949d444849">arm_compute::test::fixed_point_arithmetic::detail::max</a></div><div class="ttdeci">fixed_point&lt; T &gt; max(fixed_point&lt; T &gt; x, fixed_point&lt; T &gt; y)</div><div class="ttdef"><b>Definition:</b> <a href="tests_2validation_2_fixed_point_8h_source.xhtml#l00889">FixedPoint.h:889</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1725</div><!-- fragment -->
1726</div>
1727</div>
1728<h2 class="groupheader">Variable Documentation</h2>
1729<a class="anchor" id="a0712735973f172ac9efc7d48a31e47ad"></a>
1730<div class="memitem">
1731<div class="memproto">
1732 <table class="memname">
1733 <tr>
1734 <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1735 </tr>
1736 </table>
1737</div><div class="memdoc">
1738
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1739<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]1740
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1741<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>
1742
1743</div>
1744</div>
1745<a class="anchor" id="a4884a666a1e93fbf8c27bd7d2da3c8bb"></a>
1746<div class="memitem">
1747<div class="memproto">
1748 <table class="memname">
1749 <tr>
1750 <td class="memname">__constant uint4 idx4 = (uint4)(0, 1, 2, 3)</td>
1751 </tr>
1752 </table>
1753</div><div class="memdoc">
1754
1755<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>
1756
1757<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00500">softmax_layer_max_shift_exp_sum_parallel()</a>.</p>
1758
1759</div>
1760</div>
1761<a class="anchor" id="aa1dd94b8d98f1c6d790bdf0fc5de29e9"></a>
1762<div class="memitem">
1763<div class="memproto">
1764 <table class="memname">
1765 <tr>
1766 <td class="memname">__constant uint16 idx__ = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1767 </tr>
1768 </table>
1769</div><div class="memdoc">
1770
1771<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>
1772
1773<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]1774
1775</div>
1776</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1777<a class="anchor" id="a538b4b63f40e7b12891774e03a4f0dec"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1778<div class="memitem">
1779<div class="memproto">
1780 <table class="memname">
1781 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1782 <td class="memname">__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( -FLT_MAX )</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1783 </tr>
1784 </table>
1785</div><div class="memdoc">
1786
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1787<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]1788
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1789<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00114">softmax_layer_max()</a>.</p>
1790
1791</div>
1792</div>
1793<a class="anchor" id="af7a08044d0e491a0ee1520a24a107a2b"></a>
1794<div class="memitem">
1795<div class="memproto">
1796 <table class="memname">
1797 <tr>
1798 <td class="memname">__constant DATA_TYPE16 type_min_ = ( DATA_TYPE16 )( -FLT_MAX )</td>
1799 </tr>
1800 </table>
1801</div><div class="memdoc">
1802
1803<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>
1804
1805<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]1806
1807</div>
1808</div>
1809</div><!-- contents -->
1810</div><!-- doc-content -->
1811<!-- start footer part -->
1812<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
1813 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1814 <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>
1815 <li class="footer">Generated on Thu Dec 14 2017 23:48:34 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1816 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]1817 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1818 </ul>
1819</div>
1820</body>
1821</html>