Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 52ba29e936b8e711e8acdfe819e36f884d4f3fe1 / . / documentation / softmax__layer_8cl.xhtml
blob: 277c343cc1d9eca52b0b559dc1903cea92bc708b [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
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]43 &#160;<span id="projectnumber">18.08</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>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[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, size)&#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>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[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, size)&#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>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[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, size)&#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>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[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, size)&#160;&#160;&#160;((x) * (y))</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]137<tr class="separator:a22303c4047ec5027c1538d53964b9d0d"><td class="memSeparator" colspan="2">&#160;</td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]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, size)&#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>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[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, size)&#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: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]156<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]157<tr class="separator:ac4247ac0991e85965b7ded764e78f12c"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]158<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>
159<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>
160<tr class="separator:a9d74601bcabbb2f14bcf52385ad666dc"><td class="memSeparator" colspan="2">&#160;</td></tr>
161<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>
162<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>
163<tr class="separator:a82bc360279fcf7cf8033dca6022f21f5"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]164</table><table class="memberdecls">
165<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="var-members"></a>
166Variables</h2></td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]167<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>
168<tr class="separator:af7a08044d0e491a0ee1520a24a107a2b"><td class="memSeparator" colspan="2">&#160;</td></tr>
169<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>
170<tr class="separator:aa1dd94b8d98f1c6d790bdf0fc5de29e9"><td class="memSeparator" colspan="2">&#160;</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]171<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>
172<tr class="separator:a538b4b63f40e7b12891774e03a4f0dec"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]173<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>
174<tr class="separator:a0712735973f172ac9efc7d48a31e47ad"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]175<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>
176<tr class="separator:a4884a666a1e93fbf8c27bd7d2da3c8bb"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]177</table>
178<h2 class="groupheader">Macro Definition Documentation</h2>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]179<a class="anchor" id="a44206a4e5783c7aabacec88aad878c88"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]180<div class="memitem">
181<div class="memproto">
182 <table class="memname">
183 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]184 <td class="memname">#define ADD_OP</td>
185 <td>(</td>
186 <td class="paramtype">&#160;</td>
187 <td class="paramname">x, </td>
188 </tr>
189 <tr>
190 <td class="paramkey"></td>
191 <td></td>
192 <td class="paramtype">&#160;</td>
193 <td class="paramname">y, </td>
194 </tr>
195 <tr>
196 <td class="paramkey"></td>
197 <td></td>
198 <td class="paramtype">&#160;</td>
199 <td class="paramname">type, </td>
200 </tr>
201 <tr>
202 <td class="paramkey"></td>
203 <td></td>
204 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]205 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]206 </tr>
207 <tr>
208 <td></td>
209 <td>)</td>
210 <td></td><td>&#160;&#160;&#160;((x) + (y))</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]211 </tr>
212 </table>
213</div><div class="memdoc">
214
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]215<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00027">27</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]216
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]217<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]218
219</div>
220</div>
221<a class="anchor" id="a8cde99b1ce0f3c1dacd49261b0cf03d8"></a>
222<div class="memitem">
223<div class="memproto">
224 <table class="memname">
225 <tr>
226 <td class="memname">#define DIV_OP</td>
227 <td>(</td>
228 <td class="paramtype">&#160;</td>
229 <td class="paramname">x, </td>
230 </tr>
231 <tr>
232 <td class="paramkey"></td>
233 <td></td>
234 <td class="paramtype">&#160;</td>
235 <td class="paramname">y, </td>
236 </tr>
237 <tr>
238 <td class="paramkey"></td>
239 <td></td>
240 <td class="paramtype">&#160;</td>
241 <td class="paramname">type, </td>
242 </tr>
243 <tr>
244 <td class="paramkey"></td>
245 <td></td>
246 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]247 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]248 </tr>
249 <tr>
250 <td></td>
251 <td>)</td>
252 <td></td><td>&#160;&#160;&#160;((x) / (y))</td>
253 </tr>
254 </table>
255</div><div class="memdoc">
256
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]257<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00030">30</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]258
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]259<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00100">softmax_layer_norm()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]260
261</div>
262</div>
263<a class="anchor" id="a93cf800667317d96574477b9f0a75234"></a>
264<div class="memitem">
265<div class="memproto">
266 <table class="memname">
267 <tr>
268 <td class="memname">#define EXP_OP</td>
269 <td>(</td>
270 <td class="paramtype">&#160;</td>
271 <td class="paramname">x, </td>
272 </tr>
273 <tr>
274 <td class="paramkey"></td>
275 <td></td>
276 <td class="paramtype">&#160;</td>
277 <td class="paramname">type, </td>
278 </tr>
279 <tr>
280 <td class="paramkey"></td>
281 <td></td>
282 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]283 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]284 </tr>
285 <tr>
286 <td></td>
287 <td>)</td>
288 <td></td><td>&#160;&#160;&#160;exp((x))</td>
289 </tr>
290 </table>
291</div><div class="memdoc">
292
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]293<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00031">31</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]294
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]295<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]296
297</div>
298</div>
299<a class="anchor" id="a08246606c233e7785a497c09672f366f"></a>
300<div class="memitem">
301<div class="memproto">
302 <table class="memname">
303 <tr>
304 <td class="memname">#define GRID_SIZE&#160;&#160;&#160;1</td>
305 </tr>
306 </table>
307</div><div class="memdoc">
308
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]309<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">43</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]310
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]311<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]312
313</div>
314</div>
315<a class="anchor" id="a372393c380805985b813dbb16d589a64"></a>
316<div class="memitem">
317<div class="memproto">
318 <table class="memname">
319 <tr>
320 <td class="memname">#define LOG_VECTOR_SIZE&#160;&#160;&#160;4</td>
321 </tr>
322 </table>
323</div><div class="memdoc">
324
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]325<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00061">61</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]326
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]327<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]328
329</div>
330</div>
331<a class="anchor" id="abaa48ad818c44e415fd3f9dd0f27bf01"></a>
332<div class="memitem">
333<div class="memproto">
334 <table class="memname">
335 <tr>
336 <td class="memname">#define MAX_OP</td>
337 <td>(</td>
338 <td class="paramtype">&#160;</td>
339 <td class="paramname">x, </td>
340 </tr>
341 <tr>
342 <td class="paramkey"></td>
343 <td></td>
344 <td class="paramtype">&#160;</td>
345 <td class="paramname">y, </td>
346 </tr>
347 <tr>
348 <td class="paramkey"></td>
349 <td></td>
350 <td class="paramtype">&#160;</td>
351 <td class="paramname">type, </td>
352 </tr>
353 <tr>
354 <td class="paramkey"></td>
355 <td></td>
356 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]357 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]358 </tr>
359 <tr>
360 <td></td>
361 <td>)</td>
362 <td></td><td>&#160;&#160;&#160;max((x), (y))</td>
363 </tr>
364 </table>
365</div><div class="memdoc">
366
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]367<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00026">26</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]368
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]369<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]370
371</div>
372</div>
373<a class="anchor" id="a80b22c555ddadb47cc6ca338a9c49126"></a>
374<div class="memitem">
375<div class="memproto">
376 <table class="memname">
377 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]378 <td class="memname">#define MINVAL&#160;&#160;&#160;-FLT_MAX</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]379 </tr>
380 </table>
381</div><div class="memdoc">
382
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]383<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00037">37</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]384
385</div>
386</div>
387<a class="anchor" id="a22303c4047ec5027c1538d53964b9d0d"></a>
388<div class="memitem">
389<div class="memproto">
390 <table class="memname">
391 <tr>
392 <td class="memname">#define MUL_OP</td>
393 <td>(</td>
394 <td class="paramtype">&#160;</td>
395 <td class="paramname">x, </td>
396 </tr>
397 <tr>
398 <td class="paramkey"></td>
399 <td></td>
400 <td class="paramtype">&#160;</td>
401 <td class="paramname">y, </td>
402 </tr>
403 <tr>
404 <td class="paramkey"></td>
405 <td></td>
406 <td class="paramtype">&#160;</td>
407 <td class="paramname">type, </td>
408 </tr>
409 <tr>
410 <td class="paramkey"></td>
411 <td></td>
412 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]413 <td class="paramname">size&#160;</td>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]414 </tr>
415 <tr>
416 <td></td>
417 <td>)</td>
418 <td></td><td>&#160;&#160;&#160;((x) * (y))</td>
419 </tr>
420 </table>
421</div><div class="memdoc">
422
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]423<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00029">29</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]424
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]425<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]426
427</div>
428</div>
429<a class="anchor" id="af5b2e33e3c5fcaab3a213f26c2300170"></a>
430<div class="memitem">
431<div class="memproto">
432 <table class="memname">
433 <tr>
434 <td class="memname">#define SELECT_DATA_TYPE&#160;&#160;&#160;int</td>
435 </tr>
436 </table>
437</div><div class="memdoc">
438
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]439<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00038">38</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]440
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]441<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]442
443</div>
444</div>
445<a class="anchor" id="ac3af2d18008cbbf7247ae48fcd6e0c4e"></a>
446<div class="memitem">
447<div class="memproto">
448 <table class="memname">
449 <tr>
450 <td class="memname">#define SUB_OP</td>
451 <td>(</td>
452 <td class="paramtype">&#160;</td>
453 <td class="paramname">x, </td>
454 </tr>
455 <tr>
456 <td class="paramkey"></td>
457 <td></td>
458 <td class="paramtype">&#160;</td>
459 <td class="paramname">y, </td>
460 </tr>
461 <tr>
462 <td class="paramkey"></td>
463 <td></td>
464 <td class="paramtype">&#160;</td>
465 <td class="paramname">type, </td>
466 </tr>
467 <tr>
468 <td class="paramkey"></td>
469 <td></td>
470 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]471 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]472 </tr>
473 <tr>
474 <td></td>
475 <td>)</td>
476 <td></td><td>&#160;&#160;&#160;((x) - (y))</td>
477 </tr>
478 </table>
479</div><div class="memdoc">
480
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]481<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00028">28</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]482
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]483<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]484
485</div>
486</div>
487<a class="anchor" id="a7c78836761fa3b5b124efea237dac70f"></a>
488<div class="memitem">
489<div class="memproto">
490 <table class="memname">
491 <tr>
492 <td class="memname">#define VECTOR_SIZE&#160;&#160;&#160;16</td>
493 </tr>
494 </table>
495</div><div class="memdoc">
496
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]497<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00060">60</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]498
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]499<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]500
501</div>
502</div>
503<h2 class="groupheader">Function Documentation</h2>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]504<a class="anchor" id="a82bc360279fcf7cf8033dca6022f21f5"></a>
505<div class="memitem">
506<div class="memproto">
507 <table class="memname">
508 <tr>
509 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_parallel </td>
510 <td>(</td>
511 <td class="paramtype">__global uchar *&#160;</td>
512 <td class="paramname"><em>src_ptr</em>, </td>
513 </tr>
514 <tr>
515 <td class="paramkey"></td>
516 <td></td>
517 <td class="paramtype">uint&#160;</td>
518 <td class="paramname"><em>src_stride_x</em>, </td>
519 </tr>
520 <tr>
521 <td class="paramkey"></td>
522 <td></td>
523 <td class="paramtype">uint&#160;</td>
524 <td class="paramname"><em>src_step_x</em>, </td>
525 </tr>
526 <tr>
527 <td class="paramkey"></td>
528 <td></td>
529 <td class="paramtype">uint&#160;</td>
530 <td class="paramname"><em>src_stride_y</em>, </td>
531 </tr>
532 <tr>
533 <td class="paramkey"></td>
534 <td></td>
535 <td class="paramtype">uint&#160;</td>
536 <td class="paramname"><em>src_step_y</em>, </td>
537 </tr>
538 <tr>
539 <td class="paramkey"></td>
540 <td></td>
541 <td class="paramtype">uint&#160;</td>
542 <td class="paramname"><em>src_stride_z</em>, </td>
543 </tr>
544 <tr>
545 <td class="paramkey"></td>
546 <td></td>
547 <td class="paramtype">uint&#160;</td>
548 <td class="paramname"><em>src_step_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_offset_first_element_in_bytes</em>, </td>
555 </tr>
556 <tr>
557 <td class="paramkey"></td>
558 <td></td>
559 <td class="paramtype">__global uchar *&#160;</td>
560 <td class="paramname"><em>maxo_ptr</em>, </td>
561 </tr>
562 <tr>
563 <td class="paramkey"></td>
564 <td></td>
565 <td class="paramtype">uint&#160;</td>
566 <td class="paramname"><em>maxo_stride_x</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>maxo_step_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>maxo_stride_y</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>maxo_step_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>maxo_stride_z</em>, </td>
591 </tr>
592 <tr>
593 <td class="paramkey"></td>
594 <td></td>
595 <td class="paramtype">uint&#160;</td>
596 <td class="paramname"><em>maxo_step_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>maxo_offset_first_element_in_bytes</em>, </td>
603 </tr>
604 <tr>
605 <td class="paramkey"></td>
606 <td></td>
607 <td class="paramtype">__global uchar *&#160;</td>
608 <td class="paramname"><em>dst_ptr</em>, </td>
609 </tr>
610 <tr>
611 <td class="paramkey"></td>
612 <td></td>
613 <td class="paramtype">uint&#160;</td>
614 <td class="paramname"><em>dst_stride_x</em>, </td>
615 </tr>
616 <tr>
617 <td class="paramkey"></td>
618 <td></td>
619 <td class="paramtype">uint&#160;</td>
620 <td class="paramname"><em>dst_step_x</em>, </td>
621 </tr>
622 <tr>
623 <td class="paramkey"></td>
624 <td></td>
625 <td class="paramtype">uint&#160;</td>
626 <td class="paramname"><em>dst_stride_y</em>, </td>
627 </tr>
628 <tr>
629 <td class="paramkey"></td>
630 <td></td>
631 <td class="paramtype">uint&#160;</td>
632 <td class="paramname"><em>dst_step_y</em>, </td>
633 </tr>
634 <tr>
635 <td class="paramkey"></td>
636 <td></td>
637 <td class="paramtype">uint&#160;</td>
638 <td class="paramname"><em>dst_stride_z</em>, </td>
639 </tr>
640 <tr>
641 <td class="paramkey"></td>
642 <td></td>
643 <td class="paramtype">uint&#160;</td>
644 <td class="paramname"><em>dst_step_z</em>, </td>
645 </tr>
646 <tr>
647 <td class="paramkey"></td>
648 <td></td>
649 <td class="paramtype">uint&#160;</td>
650 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
651 </tr>
652 <tr>
653 <td class="paramkey"></td>
654 <td></td>
655 <td class="paramtype">__global uchar *&#160;</td>
656 <td class="paramname"><em>sum_ptr</em>, </td>
657 </tr>
658 <tr>
659 <td class="paramkey"></td>
660 <td></td>
661 <td class="paramtype">uint&#160;</td>
662 <td class="paramname"><em>sum_stride_x</em>, </td>
663 </tr>
664 <tr>
665 <td class="paramkey"></td>
666 <td></td>
667 <td class="paramtype">uint&#160;</td>
668 <td class="paramname"><em>sum_step_x</em>, </td>
669 </tr>
670 <tr>
671 <td class="paramkey"></td>
672 <td></td>
673 <td class="paramtype">uint&#160;</td>
674 <td class="paramname"><em>sum_stride_y</em>, </td>
675 </tr>
676 <tr>
677 <td class="paramkey"></td>
678 <td></td>
679 <td class="paramtype">uint&#160;</td>
680 <td class="paramname"><em>sum_step_y</em>, </td>
681 </tr>
682 <tr>
683 <td class="paramkey"></td>
684 <td></td>
685 <td class="paramtype">uint&#160;</td>
686 <td class="paramname"><em>sum_stride_z</em>, </td>
687 </tr>
688 <tr>
689 <td class="paramkey"></td>
690 <td></td>
691 <td class="paramtype">uint&#160;</td>
692 <td class="paramname"><em>sum_step_z</em>, </td>
693 </tr>
694 <tr>
695 <td class="paramkey"></td>
696 <td></td>
697 <td class="paramtype">uint&#160;</td>
698 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
699 </tr>
700 <tr>
701 <td class="paramkey"></td>
702 <td></td>
703 <td class="paramtype">uint&#160;</td>
704 <td class="paramname"><em>width</em>&#160;</td>
705 </tr>
706 <tr>
707 <td></td>
708 <td>)</td>
709 <td></td><td></td>
710 </tr>
711 </table>
712</div><div class="memdoc">
713
714<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>
715<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>
716<dd>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]717In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
718<dd>
719Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
720<dl class="params"><dt>Parameters</dt><dd>
721 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]722 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16/F32 </td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]723 <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>
724 <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>
725 <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>
726 <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>
727 <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>
728 <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>
729 <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>
730 <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>
731 <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>
732 <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>
733 <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>
734 <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>
735 <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>
736 <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>
737 <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>
738 <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>
739 <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>
740 <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>
741 <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>
742 <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>
743 <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>
744 <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>
745 <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>
746 <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>
747 <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>
748 <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>
749 <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>
750 <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>
751 <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>
752 <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>
753 <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>
754 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
755 </table>
756 </dd>
757</dl>
758
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]759<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">306</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]760
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]761<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00027">ADD_OP</a>, <a class="el" href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00087">arm_compute::test::validation::beta</a>, <a class="el" href="helpers_8h_source.xhtml#l00062">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00124">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="_c_l_2_arithmetic_division_8cpp_source.xhtml#l00100">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00031">EXP_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">GRID_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00069">idx4</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00026">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00029">MUL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00311">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00152">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00038">SELECT_DATA_TYPE</a>, <a class="el" href="validation_2_c_l_2_convolution_8cpp_source.xhtml#l00147">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00028">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">type_min_</a>, <a class="el" href="helpers_8h_source.xhtml#l00056">VEC_DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00050">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00053">VSTORE</a>.</p>
762<div class="fragment"><div class="line"><a name="l00312"></a><span class="lineno"> 312</span>&#160;{</div><div class="line"><a name="l00313"></a><span class="lineno"> 313</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00314"></a><span class="lineno"> 314</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="l00315"></a><span class="lineno"> 315</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="l00316"></a><span class="lineno"> 316</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="l00317"></a><span class="lineno"> 317</span>&#160;</div><div class="line"><a name="l00318"></a><span class="lineno"> 318</span>&#160; <span class="keyword">const</span> uint lid = get_local_id(0);</div><div class="line"><a name="l00319"></a><span class="lineno"> 319</span>&#160;</div><div class="line"><a name="l00320"></a><span class="lineno"> 320</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00321"></a><span class="lineno"> 321</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00322"></a><span class="lineno"> 322</span>&#160; <a class="code" href="helpers_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="l00323"></a><span class="lineno"> 323</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a> = (<a class="code" href="helpers_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="l00324"></a><span class="lineno"> 324</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00325"></a><span class="lineno"> 325</span>&#160;</div><div class="line"><a name="l00326"></a><span class="lineno"> 326</span>&#160; <span class="comment">// Define one temporary vector per work-item.</span></div><div class="line"><a name="l00327"></a><span class="lineno"> 327</span>&#160; __local <a class="code" href="helpers_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="l00328"></a><span class="lineno"> 328</span>&#160; __local DATA_TYPE max_local;</div><div class="line"><a name="l00329"></a><span class="lineno"> 329</span>&#160;</div><div class="line"><a name="l00330"></a><span class="lineno"> 330</span>&#160; __constant <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4) type_min4 = (<a class="code" href="helpers_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="l00331"></a><span class="lineno"> 331</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00332"></a><span class="lineno"> 332</span>&#160; max_val_vec = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4))type_min4;</div><div class="line"><a name="l00333"></a><span class="lineno"> 333</span>&#160; <span class="comment">// Number of elements per work-item.</span></div><div class="line"><a name="l00334"></a><span class="lineno"> 334</span>&#160; const uint row = width / GRID_SIZE;</div><div class="line"><a name="l00335"></a><span class="lineno"> 335</span>&#160; <span class="comment">// Number of iterations per work-item.</span></div><div class="line"><a name="l00336"></a><span class="lineno"> 336</span>&#160; const uint width_ = row &gt;&gt; 2;</div><div class="line"><a name="l00337"></a><span class="lineno"> 337</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00338"></a><span class="lineno"> 338</span>&#160; uint i = 0;</div><div class="line"><a name="l00339"></a><span class="lineno"> 339</span>&#160; for(; i &lt; width_; i++)</div><div class="line"><a name="l00340"></a><span class="lineno"> 340</span>&#160; {</div><div class="line"><a name="l00341"></a><span class="lineno"> 341</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00342"></a><span class="lineno"> 342</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="l00343"></a><span class="lineno"> 343</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="l00344"></a><span class="lineno"> 344</span>&#160; }</div><div class="line"><a name="l00345"></a><span class="lineno"> 345</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00346"></a><span class="lineno"> 346</span>&#160; <span class="comment">// How many work-items needed to complete the computation.</span></div><div class="line"><a name="l00347"></a><span class="lineno"> 347</span>&#160; <span class="keywordtype">int</span> boundary_workitems = (width % (GRID_SIZE * 4)) / 4;</div><div class="line"><a name="l00348"></a><span class="lineno"> 348</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00349"></a><span class="lineno"> 349</span>&#160; {</div><div class="line"><a name="l00350"></a><span class="lineno"> 350</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00351"></a><span class="lineno"> 351</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="l00352"></a><span class="lineno"> 352</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="l00353"></a><span class="lineno"> 353</span>&#160; }</div><div class="line"><a name="l00354"></a><span class="lineno"> 354</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00355"></a><span class="lineno"> 355</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</div><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; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>&#160; i--;</div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span>&#160; }</div><div class="line"><a name="l00360"></a><span class="lineno"> 360</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</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="comment">// Handle non multiple of 4</span></div><div class="line"><a name="l00363"></a><span class="lineno"> 363</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00364"></a><span class="lineno"> 364</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="l00365"></a><span class="lineno"> 365</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4)</div><div class="line"><a name="l00366"></a><span class="lineno"> 366</span>&#160; widx = <a class="code" href="helpers_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="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4));</div><div class="line"><a name="l00367"></a><span class="lineno"> 367</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="l00368"></a><span class="lineno"> 368</span>&#160; }</div><div class="line"><a name="l00369"></a><span class="lineno"> 369</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="l00370"></a><span class="lineno"> 370</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="l00371"></a><span class="lineno"> 371</span>&#160; tmp_local[lid] = max_val_vec;</div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>&#160;</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span>&#160;</div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 256)</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>&#160; {</div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>&#160; <span class="keywordflow">if</span>(lid &lt; 128)</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>&#160; {</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</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="l00380"></a><span class="lineno"> 380</span>&#160; }</div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>&#160; }</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 128)</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span>&#160; {</div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>&#160; <span class="keywordflow">if</span>(lid &lt; 64)</div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span>&#160; {</div><div class="line"><a name="l00387"></a><span class="lineno"> 387</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="l00388"></a><span class="lineno"> 388</span>&#160; }</div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00390"></a><span class="lineno"> 390</span>&#160; }</div><div class="line"><a name="l00391"></a><span class="lineno"> 391</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 64)</div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span>&#160; {</div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span>&#160; <span class="keywordflow">if</span>(lid &lt; 32)</div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span>&#160; {</div><div class="line"><a name="l00395"></a><span class="lineno"> 395</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="l00396"></a><span class="lineno"> 396</span>&#160; }</div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>&#160; }</div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 32)</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span>&#160; {</div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>&#160; <span class="keywordflow">if</span>(lid &lt; 16)</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; 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="l00404"></a><span class="lineno"> 404</span>&#160; }</div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00406"></a><span class="lineno"> 406</span>&#160; }</div><div class="line"><a name="l00407"></a><span class="lineno"> 407</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 16)</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span>&#160; {</div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>&#160; <span class="keywordflow">if</span>(lid &lt; 8)</div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span>&#160; {</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</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="l00412"></a><span class="lineno"> 412</span>&#160; }</div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 8)</div><div class="line"><a name="l00416"></a><span class="lineno"> 416</span>&#160; {</div><div class="line"><a name="l00417"></a><span class="lineno"> 417</span>&#160; <span class="keywordflow">if</span>(lid &lt; 4)</div><div class="line"><a name="l00418"></a><span class="lineno"> 418</span>&#160; {</div><div class="line"><a name="l00419"></a><span class="lineno"> 419</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="l00420"></a><span class="lineno"> 420</span>&#160; }</div><div class="line"><a name="l00421"></a><span class="lineno"> 421</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00422"></a><span class="lineno"> 422</span>&#160; }</div><div class="line"><a name="l00423"></a><span class="lineno"> 423</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 4)</div><div class="line"><a name="l00424"></a><span class="lineno"> 424</span>&#160; {</div><div class="line"><a name="l00425"></a><span class="lineno"> 425</span>&#160; <span class="keywordflow">if</span>(lid &lt; 2)</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; 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="l00428"></a><span class="lineno"> 428</span>&#160; }</div><div class="line"><a name="l00429"></a><span class="lineno"> 429</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00430"></a><span class="lineno"> 430</span>&#160; }</div><div class="line"><a name="l00431"></a><span class="lineno"> 431</span>&#160; <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00432"></a><span class="lineno"> 432</span>&#160; {</div><div class="line"><a name="l00433"></a><span class="lineno"> 433</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="l00434"></a><span class="lineno"> 434</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="l00435"></a><span class="lineno"> 435</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="l00436"></a><span class="lineno"> 436</span>&#160; max_local = max_val_vec.s0;</div><div class="line"><a name="l00437"></a><span class="lineno"> 437</span>&#160; }</div><div class="line"><a name="l00438"></a><span class="lineno"> 438</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00439"></a><span class="lineno"> 439</span>&#160;</div><div class="line"><a name="l00440"></a><span class="lineno"> 440</span>&#160; <span class="comment">/* Second section */</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">// Set sum vector</span></div><div class="line"><a name="l00443"></a><span class="lineno"> 443</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00444"></a><span class="lineno"> 444</span>&#160; sum1D = 0;</div><div class="line"><a name="l00445"></a><span class="lineno"> 445</span>&#160; DATA_TYPE max_val = max_local;</div><div class="line"><a name="l00446"></a><span class="lineno"> 446</span>&#160;</div><div class="line"><a name="l00447"></a><span class="lineno"> 447</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00448"></a><span class="lineno"> 448</span>&#160; for(i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00449"></a><span class="lineno"> 449</span>&#160; {</div><div class="line"><a name="l00450"></a><span class="lineno"> 450</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00451"></a><span class="lineno"> 451</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="l00452"></a><span class="lineno"> 452</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="l00453"></a><span class="lineno"> 453</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00454"></a><span class="lineno"> 454</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="l00455"></a><span class="lineno"> 455</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00456"></a><span class="lineno"> 456</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="l00457"></a><span class="lineno"> 457</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00458"></a><span class="lineno"> 458</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="l00459"></a><span class="lineno"> 459</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="l00460"></a><span class="lineno"> 460</span>&#160; }</div><div class="line"><a name="l00461"></a><span class="lineno"> 461</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00462"></a><span class="lineno"> 462</span>&#160; boundary_workitems = (width % (GRID_SIZE * 4)) / 4;</div><div class="line"><a name="l00463"></a><span class="lineno"> 463</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00464"></a><span class="lineno"> 464</span>&#160; {</div><div class="line"><a name="l00465"></a><span class="lineno"> 465</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00466"></a><span class="lineno"> 466</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="l00467"></a><span class="lineno"> 467</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="l00468"></a><span class="lineno"> 468</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00469"></a><span class="lineno"> 469</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="l00470"></a><span class="lineno"> 470</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00471"></a><span class="lineno"> 471</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="l00472"></a><span class="lineno"> 472</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00473"></a><span class="lineno"> 473</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="l00474"></a><span class="lineno"> 474</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="l00475"></a><span class="lineno"> 475</span>&#160; }</div><div class="line"><a name="l00476"></a><span class="lineno"> 476</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00477"></a><span class="lineno"> 477</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00478"></a><span class="lineno"> 478</span>&#160; {</div><div class="line"><a name="l00479"></a><span class="lineno"> 479</span>&#160; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00480"></a><span class="lineno"> 480</span>&#160; i--;</div><div class="line"><a name="l00481"></a><span class="lineno"> 481</span>&#160; }</div><div class="line"><a name="l00482"></a><span class="lineno"> 482</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00483"></a><span class="lineno"> 483</span>&#160; {</div><div class="line"><a name="l00484"></a><span class="lineno"> 484</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="l00485"></a><span class="lineno"> 485</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00486"></a><span class="lineno"> 486</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="l00487"></a><span class="lineno"> 487</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="l00488"></a><span class="lineno"> 488</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00489"></a><span class="lineno"> 489</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="l00490"></a><span class="lineno"> 490</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00491"></a><span class="lineno"> 491</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="l00492"></a><span class="lineno"> 492</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4)</div><div class="line"><a name="l00493"></a><span class="lineno"> 493</span>&#160; widx = <a class="code" href="helpers_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="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4));</div><div class="line"><a name="l00494"></a><span class="lineno"> 494</span>&#160; data = select(0, data, widx);</div><div class="line"><a name="l00495"></a><span class="lineno"> 495</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00496"></a><span class="lineno"> 496</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="l00497"></a><span class="lineno"> 497</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="l00498"></a><span class="lineno"> 498</span>&#160; }</div><div class="line"><a name="l00499"></a><span class="lineno"> 499</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="l00500"></a><span class="lineno"> 500</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="l00501"></a><span class="lineno"> 501</span>&#160; tmp_local[lid] = sum1D;</div><div class="line"><a name="l00502"></a><span class="lineno"> 502</span>&#160;</div><div class="line"><a name="l00503"></a><span class="lineno"> 503</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00504"></a><span class="lineno"> 504</span>&#160;</div><div class="line"><a name="l00505"></a><span class="lineno"> 505</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 256)</div><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; <span class="keywordflow">if</span>(lid &lt; 128)</div><div class="line"><a name="l00508"></a><span class="lineno"> 508</span>&#160; {</div><div class="line"><a name="l00509"></a><span class="lineno"> 509</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="l00510"></a><span class="lineno"> 510</span>&#160; }</div><div class="line"><a name="l00511"></a><span class="lineno"> 511</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00512"></a><span class="lineno"> 512</span>&#160; }</div><div class="line"><a name="l00513"></a><span class="lineno"> 513</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 128)</div><div class="line"><a name="l00514"></a><span class="lineno"> 514</span>&#160; {</div><div class="line"><a name="l00515"></a><span class="lineno"> 515</span>&#160; <span class="keywordflow">if</span>(lid &lt; 64)</div><div class="line"><a name="l00516"></a><span class="lineno"> 516</span>&#160; {</div><div class="line"><a name="l00517"></a><span class="lineno"> 517</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="l00518"></a><span class="lineno"> 518</span>&#160; }</div><div class="line"><a name="l00519"></a><span class="lineno"> 519</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00520"></a><span class="lineno"> 520</span>&#160; }</div><div class="line"><a name="l00521"></a><span class="lineno"> 521</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 64)</div><div class="line"><a name="l00522"></a><span class="lineno"> 522</span>&#160; {</div><div class="line"><a name="l00523"></a><span class="lineno"> 523</span>&#160; <span class="keywordflow">if</span>(lid &lt; 32)</div><div class="line"><a name="l00524"></a><span class="lineno"> 524</span>&#160; {</div><div class="line"><a name="l00525"></a><span class="lineno"> 525</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="l00526"></a><span class="lineno"> 526</span>&#160; }</div><div class="line"><a name="l00527"></a><span class="lineno"> 527</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00528"></a><span class="lineno"> 528</span>&#160; }</div><div class="line"><a name="l00529"></a><span class="lineno"> 529</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 32)</div><div class="line"><a name="l00530"></a><span class="lineno"> 530</span>&#160; {</div><div class="line"><a name="l00531"></a><span class="lineno"> 531</span>&#160; <span class="keywordflow">if</span>(lid &lt; 16)</div><div class="line"><a name="l00532"></a><span class="lineno"> 532</span>&#160; {</div><div class="line"><a name="l00533"></a><span class="lineno"> 533</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="l00534"></a><span class="lineno"> 534</span>&#160; }</div><div class="line"><a name="l00535"></a><span class="lineno"> 535</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00536"></a><span class="lineno"> 536</span>&#160; }</div><div class="line"><a name="l00537"></a><span class="lineno"> 537</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 16)</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="keywordflow">if</span>(lid &lt; 8)</div><div class="line"><a name="l00540"></a><span class="lineno"> 540</span>&#160; {</div><div class="line"><a name="l00541"></a><span class="lineno"> 541</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="l00542"></a><span class="lineno"> 542</span>&#160; }</div><div class="line"><a name="l00543"></a><span class="lineno"> 543</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00544"></a><span class="lineno"> 544</span>&#160; }</div><div class="line"><a name="l00545"></a><span class="lineno"> 545</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 8)</div><div class="line"><a name="l00546"></a><span class="lineno"> 546</span>&#160; {</div><div class="line"><a name="l00547"></a><span class="lineno"> 547</span>&#160; <span class="keywordflow">if</span>(lid &lt; 4)</div><div class="line"><a name="l00548"></a><span class="lineno"> 548</span>&#160; {</div><div class="line"><a name="l00549"></a><span class="lineno"> 549</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="l00550"></a><span class="lineno"> 550</span>&#160; }</div><div class="line"><a name="l00551"></a><span class="lineno"> 551</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00552"></a><span class="lineno"> 552</span>&#160; }</div><div class="line"><a name="l00553"></a><span class="lineno"> 553</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 4)</div><div class="line"><a name="l00554"></a><span class="lineno"> 554</span>&#160; {</div><div class="line"><a name="l00555"></a><span class="lineno"> 555</span>&#160; <span class="keywordflow">if</span>(lid &lt; 2)</div><div class="line"><a name="l00556"></a><span class="lineno"> 556</span>&#160; {</div><div class="line"><a name="l00557"></a><span class="lineno"> 557</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="l00558"></a><span class="lineno"> 558</span>&#160; }</div><div class="line"><a name="l00559"></a><span class="lineno"> 559</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00560"></a><span class="lineno"> 560</span>&#160; }</div><div class="line"><a name="l00561"></a><span class="lineno"> 561</span>&#160; <span class="keywordflow">if</span>(lid == 0)</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; 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="l00564"></a><span class="lineno"> 564</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00565"></a><span class="lineno"> 565</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="l00566"></a><span class="lineno"> 566</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="l00567"></a><span class="lineno"> 567</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</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;}</div><div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ab146b9cbab6e73e7588b240dc709fe01"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">arm_compute::test::validation::beta</a></div><div class="ttdeci">beta</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00087">GEMM.cpp:87</a></div></div>
763<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#l00050">helpers.h:50</a></div></div>
764<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#l00031">softmax_layer.cl:31</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]765<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>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]766<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#l00026">softmax_layer.cl:26</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]767<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>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]768<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#l00062">softmax_layer.cl:62</a></div></div>
769<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#l00027">softmax_layer.cl:27</a></div></div>
770<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#l00311">helpers.h:311</a></div></div>
771<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#l00053">helpers.h:53</a></div></div>
772<div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00056">helpers.h:56</a></div></div>
773<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#l00038">softmax_layer.cl:38</a></div></div>
774<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#l00069">softmax_layer.cl:69</a></div></div>
775<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_arithmetic_division_8cpp_source.xhtml#l00100">ArithmeticDivision.cpp:100</a></div></div>
776<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#l00029">softmax_layer.cl:29</a></div></div>
777<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#l00150">helpers.h:150</a></div></div>
778<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#l00152">helpers.h:152</a></div></div>
779<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#l00028">softmax_layer.cl:28</a></div></div>
780<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#l00037">softmax_layer.cl:37</a></div></div>
781<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#l00124">helpers.h:124</a></div></div>
782<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#l00043">softmax_layer.cl:43</a></div></div>
783<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
784<div class="ttc" id="helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00062">helpers.h:62</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]785</div><!-- fragment -->
786</div>
787</div>
788<a class="anchor" id="a9d74601bcabbb2f14bcf52385ad666dc"></a>
789<div class="memitem">
790<div class="memproto">
791 <table class="memname">
792 <tr>
793 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_serial </td>
794 <td>(</td>
795 <td class="paramtype">__global uchar *&#160;</td>
796 <td class="paramname"><em>src_ptr</em>, </td>
797 </tr>
798 <tr>
799 <td class="paramkey"></td>
800 <td></td>
801 <td class="paramtype">uint&#160;</td>
802 <td class="paramname"><em>src_stride_x</em>, </td>
803 </tr>
804 <tr>
805 <td class="paramkey"></td>
806 <td></td>
807 <td class="paramtype">uint&#160;</td>
808 <td class="paramname"><em>src_step_x</em>, </td>
809 </tr>
810 <tr>
811 <td class="paramkey"></td>
812 <td></td>
813 <td class="paramtype">uint&#160;</td>
814 <td class="paramname"><em>src_stride_y</em>, </td>
815 </tr>
816 <tr>
817 <td class="paramkey"></td>
818 <td></td>
819 <td class="paramtype">uint&#160;</td>
820 <td class="paramname"><em>src_step_y</em>, </td>
821 </tr>
822 <tr>
823 <td class="paramkey"></td>
824 <td></td>
825 <td class="paramtype">uint&#160;</td>
826 <td class="paramname"><em>src_stride_z</em>, </td>
827 </tr>
828 <tr>
829 <td class="paramkey"></td>
830 <td></td>
831 <td class="paramtype">uint&#160;</td>
832 <td class="paramname"><em>src_step_z</em>, </td>
833 </tr>
834 <tr>
835 <td class="paramkey"></td>
836 <td></td>
837 <td class="paramtype">uint&#160;</td>
838 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
839 </tr>
840 <tr>
841 <td class="paramkey"></td>
842 <td></td>
843 <td class="paramtype">__global uchar *&#160;</td>
844 <td class="paramname"><em>maxo_ptr</em>, </td>
845 </tr>
846 <tr>
847 <td class="paramkey"></td>
848 <td></td>
849 <td class="paramtype">uint&#160;</td>
850 <td class="paramname"><em>maxo_stride_x</em>, </td>
851 </tr>
852 <tr>
853 <td class="paramkey"></td>
854 <td></td>
855 <td class="paramtype">uint&#160;</td>
856 <td class="paramname"><em>maxo_step_x</em>, </td>
857 </tr>
858 <tr>
859 <td class="paramkey"></td>
860 <td></td>
861 <td class="paramtype">uint&#160;</td>
862 <td class="paramname"><em>maxo_stride_y</em>, </td>
863 </tr>
864 <tr>
865 <td class="paramkey"></td>
866 <td></td>
867 <td class="paramtype">uint&#160;</td>
868 <td class="paramname"><em>maxo_step_y</em>, </td>
869 </tr>
870 <tr>
871 <td class="paramkey"></td>
872 <td></td>
873 <td class="paramtype">uint&#160;</td>
874 <td class="paramname"><em>maxo_stride_z</em>, </td>
875 </tr>
876 <tr>
877 <td class="paramkey"></td>
878 <td></td>
879 <td class="paramtype">uint&#160;</td>
880 <td class="paramname"><em>maxo_step_z</em>, </td>
881 </tr>
882 <tr>
883 <td class="paramkey"></td>
884 <td></td>
885 <td class="paramtype">uint&#160;</td>
886 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
887 </tr>
888 <tr>
889 <td class="paramkey"></td>
890 <td></td>
891 <td class="paramtype">__global uchar *&#160;</td>
892 <td class="paramname"><em>dst_ptr</em>, </td>
893 </tr>
894 <tr>
895 <td class="paramkey"></td>
896 <td></td>
897 <td class="paramtype">uint&#160;</td>
898 <td class="paramname"><em>dst_stride_x</em>, </td>
899 </tr>
900 <tr>
901 <td class="paramkey"></td>
902 <td></td>
903 <td class="paramtype">uint&#160;</td>
904 <td class="paramname"><em>dst_step_x</em>, </td>
905 </tr>
906 <tr>
907 <td class="paramkey"></td>
908 <td></td>
909 <td class="paramtype">uint&#160;</td>
910 <td class="paramname"><em>dst_stride_y</em>, </td>
911 </tr>
912 <tr>
913 <td class="paramkey"></td>
914 <td></td>
915 <td class="paramtype">uint&#160;</td>
916 <td class="paramname"><em>dst_step_y</em>, </td>
917 </tr>
918 <tr>
919 <td class="paramkey"></td>
920 <td></td>
921 <td class="paramtype">uint&#160;</td>
922 <td class="paramname"><em>dst_stride_z</em>, </td>
923 </tr>
924 <tr>
925 <td class="paramkey"></td>
926 <td></td>
927 <td class="paramtype">uint&#160;</td>
928 <td class="paramname"><em>dst_step_z</em>, </td>
929 </tr>
930 <tr>
931 <td class="paramkey"></td>
932 <td></td>
933 <td class="paramtype">uint&#160;</td>
934 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
935 </tr>
936 <tr>
937 <td class="paramkey"></td>
938 <td></td>
939 <td class="paramtype">__global uchar *&#160;</td>
940 <td class="paramname"><em>sum_ptr</em>, </td>
941 </tr>
942 <tr>
943 <td class="paramkey"></td>
944 <td></td>
945 <td class="paramtype">uint&#160;</td>
946 <td class="paramname"><em>sum_stride_x</em>, </td>
947 </tr>
948 <tr>
949 <td class="paramkey"></td>
950 <td></td>
951 <td class="paramtype">uint&#160;</td>
952 <td class="paramname"><em>sum_step_x</em>, </td>
953 </tr>
954 <tr>
955 <td class="paramkey"></td>
956 <td></td>
957 <td class="paramtype">uint&#160;</td>
958 <td class="paramname"><em>sum_stride_y</em>, </td>
959 </tr>
960 <tr>
961 <td class="paramkey"></td>
962 <td></td>
963 <td class="paramtype">uint&#160;</td>
964 <td class="paramname"><em>sum_step_y</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>sum_stride_z</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>sum_step_z</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>sum_offset_first_element_in_bytes</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>width</em>&#160;</td>
989 </tr>
990 <tr>
991 <td></td>
992 <td>)</td>
993 <td></td><td></td>
994 </tr>
995 </table>
996</div><div class="memdoc">
997
998<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>
999<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>
1000<dd>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1001In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
1002<dd>
1003Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
1004<dl class="params"><dt>Parameters</dt><dd>
1005 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1006 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16/F32 </td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1007 <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>
1008 <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>
1009 <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>
1010 <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>
1011 <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>
1012 <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>
1013 <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>
1014 <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>
1015 <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>
1016 <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>
1017 <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>
1018 <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>
1019 <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>
1020 <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>
1021 <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>
1022 <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>
1023 <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>
1024 <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>
1025 <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>
1026 <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>
1027 <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>
1028 <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>
1029 <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>
1030 <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>
1031 <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>
1032 <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>
1033 <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>
1034 <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>
1035 <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>
1036 <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>
1037 <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>
1038 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1039 </table>
1040 </dd>
1041</dl>
1042
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1043<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">157</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1044
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1045<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00027">ADD_OP</a>, <a class="el" href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00087">arm_compute::test::validation::beta</a>, <a class="el" href="helpers_8h_source.xhtml#l00059">CL_VEC_DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00062">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00124">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="_c_l_2_arithmetic_division_8cpp_source.xhtml#l00100">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00031">EXP_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00045">EXPAND</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00063">idx__</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00061">LOG_VECTOR_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00026">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00029">MUL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00311">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00152">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00038">SELECT_DATA_TYPE</a>, <a class="el" href="validation_2_c_l_2_convolution_8cpp_source.xhtml#l00147">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00028">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">type_min_</a>, <a class="el" href="helpers_8h_source.xhtml#l00056">VEC_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00060">VECTOR_SIZE</a>, <a class="el" href="helpers_8h_source.xhtml#l00050">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00053">VSTORE</a>.</p>
1046<div class="fragment"><div class="line"><a name="l00163"></a><span class="lineno"> 163</span>&#160;{</div><div class="line"><a name="l00164"></a><span class="lineno"> 164</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00165"></a><span class="lineno"> 165</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="l00166"></a><span class="lineno"> 166</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="l00167"></a><span class="lineno"> 167</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="l00168"></a><span class="lineno"> 168</span>&#160;</div><div class="line"><a name="l00169"></a><span class="lineno"> 169</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00170"></a><span class="lineno"> 170</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00171"></a><span class="lineno"> 171</span>&#160; <a class="code" href="helpers_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="l00172"></a><span class="lineno"> 172</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a> = (<a class="code" href="helpers_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="l00173"></a><span class="lineno"> 173</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00174"></a><span class="lineno"> 174</span>&#160;</div><div class="line"><a name="l00175"></a><span class="lineno"> 175</span>&#160; <span class="comment">// Initialize local maximum</span></div><div class="line"><a name="l00176"></a><span class="lineno"> 176</span>&#160; <a class="code" href="helpers_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="l00177"></a><span class="lineno"> 177</span>&#160; max_val_vec = (<a class="code" href="helpers_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="l00178"></a><span class="lineno"> 178</span>&#160;</div><div class="line"><a name="l00179"></a><span class="lineno"> 179</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00180"></a><span class="lineno"> 180</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="l00181"></a><span class="lineno"> 181</span>&#160; for(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00182"></a><span class="lineno"> 182</span>&#160; {</div><div class="line"><a name="l00183"></a><span class="lineno"> 183</span>&#160; <a class="code" href="helpers_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="l00184"></a><span class="lineno"> 184</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="l00185"></a><span class="lineno"> 185</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="l00186"></a><span class="lineno"> 186</span>&#160; }</div><div class="line"><a name="l00187"></a><span class="lineno"> 187</span>&#160;</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00189"></a><span class="lineno"> 189</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00190"></a><span class="lineno"> 190</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="l00191"></a><span class="lineno"> 191</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, VECTOR_SIZE)</div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>&#160; widx = <a class="code" href="helpers_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="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, VECTOR_SIZE));</div><div class="line"><a name="l00193"></a><span class="lineno"> 193</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="l00194"></a><span class="lineno"> 194</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="l00195"></a><span class="lineno"> 195</span>&#160;</div><div class="line"><a name="l00196"></a><span class="lineno"> 196</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00198"></a><span class="lineno"> 198</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="l00199"></a><span class="lineno"> 199</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="l00200"></a><span class="lineno"> 200</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></div><div class="line"><a name="l00201"></a><span class="lineno"> 201</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="l00202"></a><span class="lineno"> 202</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="l00203"></a><span class="lineno"> 203</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00204"></a><span class="lineno"> 204</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="l00205"></a><span class="lineno"> 205</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="l00206"></a><span class="lineno"> 206</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="l00207"></a><span class="lineno"> 207</span>&#160; <span class="comment">// Store result</span></div><div class="line"><a name="l00208"></a><span class="lineno"> 208</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="l00209"></a><span class="lineno"> 209</span>&#160;</div><div class="line"><a name="l00210"></a><span class="lineno"> 210</span>&#160; <span class="comment">/* Second section */</span></div><div class="line"><a name="l00211"></a><span class="lineno"> 211</span>&#160;</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00213"></a><span class="lineno"> 213</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="l00214"></a><span class="lineno"> 214</span>&#160;</div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>&#160; <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00216"></a><span class="lineno"> 216</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>&#160; sum1D = 0;</div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>&#160;</div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span>&#160; for(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00221"></a><span class="lineno"> 221</span>&#160; {</div><div class="line"><a name="l00222"></a><span class="lineno"> 222</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00223"></a><span class="lineno"> 223</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="l00224"></a><span class="lineno"> 224</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="l00225"></a><span class="lineno"> 225</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00226"></a><span class="lineno"> 226</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="l00227"></a><span class="lineno"> 227</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00228"></a><span class="lineno"> 228</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="l00229"></a><span class="lineno"> 229</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VECTOR_SIZE)</div><div class="line"><a name="l00230"></a><span class="lineno"> 230</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="l00231"></a><span class="lineno"> 231</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="l00232"></a><span class="lineno"> 232</span>&#160; }</div><div class="line"><a name="l00233"></a><span class="lineno"> 233</span>&#160;</div><div class="line"><a name="l00234"></a><span class="lineno"> 234</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00235"></a><span class="lineno"> 235</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00236"></a><span class="lineno"> 236</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="l00237"></a><span class="lineno"> 237</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="l00238"></a><span class="lineno"> 238</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00239"></a><span class="lineno"> 239</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="l00240"></a><span class="lineno"> 240</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00241"></a><span class="lineno"> 241</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="l00242"></a><span class="lineno"> 242</span>&#160; widx = <a class="code" href="helpers_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="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, VECTOR_SIZE));</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span>&#160; data = select(0, data, widx);</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VECTOR_SIZE)</div><div class="line"><a name="l00245"></a><span class="lineno"> 245</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="l00246"></a><span class="lineno"> 246</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="l00247"></a><span class="lineno"> 247</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="l00248"></a><span class="lineno"> 248</span>&#160;</div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>&#160; <span class="comment">// Perform sum reduction</span></div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00251"></a><span class="lineno"> 251</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="l00252"></a><span class="lineno"> 252</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="l00253"></a><span class="lineno"> 253</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></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;<span class="preprocessor">#endif </span><span class="comment">/* VECTOR SIZE 8 END */</span><span class="preprocessor"></span></div><div class="line"><a name="l00256"></a><span class="lineno"> 256</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00257"></a><span class="lineno"> 257</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="l00258"></a><span class="lineno"> 258</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="l00259"></a><span class="lineno"> 259</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="l00260"></a><span class="lineno"> 260</span>&#160;</div><div class="line"><a name="l00261"></a><span class="lineno"> 261</span>&#160; <span class="comment">// Calculate and store result</span></div><div class="line"><a name="l00262"></a><span class="lineno"> 262</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div><div class="line"><a name="l00263"></a><span class="lineno"> 263</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#l00063">softmax_layer.cl:63</a></div></div>
1047<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ab146b9cbab6e73e7588b240dc709fe01"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">arm_compute::test::validation::beta</a></div><div class="ttdeci">beta</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00087">GEMM.cpp:87</a></div></div>
1048<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#l00050">helpers.h:50</a></div></div>
1049<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#l00031">softmax_layer.cl:31</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1050<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>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1051<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#l00061">softmax_layer.cl:61</a></div></div>
1052<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#l00026">softmax_layer.cl:26</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1053<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>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1054<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#l00062">softmax_layer.cl:62</a></div></div>
1055<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#l00027">softmax_layer.cl:27</a></div></div>
1056<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#l00311">helpers.h:311</a></div></div>
1057<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#l00053">helpers.h:53</a></div></div>
1058<div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00056">helpers.h:56</a></div></div>
1059<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#l00060">softmax_layer.cl:60</a></div></div>
1060<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#l00038">softmax_layer.cl:38</a></div></div>
1061<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_arithmetic_division_8cpp_source.xhtml#l00100">ArithmeticDivision.cpp:100</a></div></div>
1062<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#l00029">softmax_layer.cl:29</a></div></div>
1063<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#l00150">helpers.h:150</a></div></div>
1064<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#l00152">helpers.h:152</a></div></div>
1065<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#l00059">helpers.h:59</a></div></div>
1066<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#l00028">softmax_layer.cl:28</a></div></div>
1067<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#l00124">helpers.h:124</a></div></div>
1068<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#l00045">helpers.h:45</a></div></div>
1069<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
1070<div class="ttc" id="helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00062">helpers.h:62</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1071</div><!-- fragment -->
1072</div>
1073</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1074<a class="anchor" id="ac4247ac0991e85965b7ded764e78f12c"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1075<div class="memitem">
1076<div class="memproto">
1077 <table class="memname">
1078 <tr>
1079 <td class="memname">__kernel void softmax_layer_norm </td>
1080 <td>(</td>
1081 <td class="paramtype">__global uchar *&#160;</td>
1082 <td class="paramname"><em>src_ptr</em>, </td>
1083 </tr>
1084 <tr>
1085 <td class="paramkey"></td>
1086 <td></td>
1087 <td class="paramtype">uint&#160;</td>
1088 <td class="paramname"><em>src_stride_x</em>, </td>
1089 </tr>
1090 <tr>
1091 <td class="paramkey"></td>
1092 <td></td>
1093 <td class="paramtype">uint&#160;</td>
1094 <td class="paramname"><em>src_step_x</em>, </td>
1095 </tr>
1096 <tr>
1097 <td class="paramkey"></td>
1098 <td></td>
1099 <td class="paramtype">uint&#160;</td>
1100 <td class="paramname"><em>src_stride_y</em>, </td>
1101 </tr>
1102 <tr>
1103 <td class="paramkey"></td>
1104 <td></td>
1105 <td class="paramtype">uint&#160;</td>
1106 <td class="paramname"><em>src_step_y</em>, </td>
1107 </tr>
1108 <tr>
1109 <td class="paramkey"></td>
1110 <td></td>
1111 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1112 <td class="paramname"><em>src_stride_z</em>, </td>
1113 </tr>
1114 <tr>
1115 <td class="paramkey"></td>
1116 <td></td>
1117 <td class="paramtype">uint&#160;</td>
1118 <td class="paramname"><em>src_step_z</em>, </td>
1119 </tr>
1120 <tr>
1121 <td class="paramkey"></td>
1122 <td></td>
1123 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1124 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1125 </tr>
1126 <tr>
1127 <td class="paramkey"></td>
1128 <td></td>
1129 <td class="paramtype">__global uchar *&#160;</td>
1130 <td class="paramname"><em>sum_ptr</em>, </td>
1131 </tr>
1132 <tr>
1133 <td class="paramkey"></td>
1134 <td></td>
1135 <td class="paramtype">uint&#160;</td>
1136 <td class="paramname"><em>sum_stride_x</em>, </td>
1137 </tr>
1138 <tr>
1139 <td class="paramkey"></td>
1140 <td></td>
1141 <td class="paramtype">uint&#160;</td>
1142 <td class="paramname"><em>sum_step_x</em>, </td>
1143 </tr>
1144 <tr>
1145 <td class="paramkey"></td>
1146 <td></td>
1147 <td class="paramtype">uint&#160;</td>
1148 <td class="paramname"><em>sum_stride_y</em>, </td>
1149 </tr>
1150 <tr>
1151 <td class="paramkey"></td>
1152 <td></td>
1153 <td class="paramtype">uint&#160;</td>
1154 <td class="paramname"><em>sum_step_y</em>, </td>
1155 </tr>
1156 <tr>
1157 <td class="paramkey"></td>
1158 <td></td>
1159 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1160 <td class="paramname"><em>sum_stride_z</em>, </td>
1161 </tr>
1162 <tr>
1163 <td class="paramkey"></td>
1164 <td></td>
1165 <td class="paramtype">uint&#160;</td>
1166 <td class="paramname"><em>sum_step_z</em>, </td>
1167 </tr>
1168 <tr>
1169 <td class="paramkey"></td>
1170 <td></td>
1171 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1172 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1173 </tr>
1174 <tr>
1175 <td class="paramkey"></td>
1176 <td></td>
1177 <td class="paramtype">__global uchar *&#160;</td>
1178 <td class="paramname"><em>dst_ptr</em>, </td>
1179 </tr>
1180 <tr>
1181 <td class="paramkey"></td>
1182 <td></td>
1183 <td class="paramtype">uint&#160;</td>
1184 <td class="paramname"><em>dst_stride_x</em>, </td>
1185 </tr>
1186 <tr>
1187 <td class="paramkey"></td>
1188 <td></td>
1189 <td class="paramtype">uint&#160;</td>
1190 <td class="paramname"><em>dst_step_x</em>, </td>
1191 </tr>
1192 <tr>
1193 <td class="paramkey"></td>
1194 <td></td>
1195 <td class="paramtype">uint&#160;</td>
1196 <td class="paramname"><em>dst_stride_y</em>, </td>
1197 </tr>
1198 <tr>
1199 <td class="paramkey"></td>
1200 <td></td>
1201 <td class="paramtype">uint&#160;</td>
1202 <td class="paramname"><em>dst_step_y</em>, </td>
1203 </tr>
1204 <tr>
1205 <td class="paramkey"></td>
1206 <td></td>
1207 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1208 <td class="paramname"><em>dst_stride_z</em>, </td>
1209 </tr>
1210 <tr>
1211 <td class="paramkey"></td>
1212 <td></td>
1213 <td class="paramtype">uint&#160;</td>
1214 <td class="paramname"><em>dst_step_z</em>, </td>
1215 </tr>
1216 <tr>
1217 <td class="paramkey"></td>
1218 <td></td>
1219 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1220 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
1221 </tr>
1222 <tr>
1223 <td></td>
1224 <td>)</td>
1225 <td></td><td></td>
1226 </tr>
1227 </table>
1228</div><div class="memdoc">
1229
1230<p>Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel. </p>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1231<dl class="section note"><dt>Note</dt><dd>Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short</dd></dl>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1232<dl class="params"><dt>Parameters</dt><dd>
1233 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1234 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1235 <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>
1236 <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>
1237 <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>
1238 <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]1239 <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>
1240 <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]1241 <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]1242 <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]1243 <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>
1244 <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>
1245 <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>
1246 <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]1247 <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>
1248 <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]1249 <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]1250 <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]1251 <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>
1252 <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>
1253 <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>
1254 <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]1255 <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>
1256 <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]1257 <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>
1258 </table>
1259 </dd>
1260</dl>
1261
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1262<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00100">100</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]1263
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1264<p>References <a class="el" href="helpers_8h_source.xhtml#l00124">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00121">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00030">DIV_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00311">offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00056">VEC_DATA_TYPE</a>.</p>
1265<div class="fragment"><div class="line"><a name="l00104"></a><span class="lineno"> 104</span>&#160;{</div><div class="line"><a name="l00105"></a><span class="lineno"> 105</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</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="l00107"></a><span class="lineno"> 107</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="l00108"></a><span class="lineno"> 108</span>&#160;</div><div class="line"><a name="l00109"></a><span class="lineno"> 109</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00110"></a><span class="lineno"> 110</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="l00111"></a><span class="lineno"> 111</span>&#160; <a class="code" href="helpers_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="l00112"></a><span class="lineno"> 112</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="l00113"></a><span class="lineno"> 113</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="l00114"></a><span class="lineno"> 114</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#l00121">helpers.h:121</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1266<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>
1267<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>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1268<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#l00030">softmax_layer.cl:30</a></div></div>
1269<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#l00311">helpers.h:311</a></div></div>
1270<div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00056">helpers.h:56</a></div></div>
1271<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_arithmetic_division_8cpp_source.xhtml#l00100">ArithmeticDivision.cpp:100</a></div></div>
1272<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#l00150">helpers.h:150</a></div></div>
1273<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#l00124">helpers.h:124</a></div></div>
1274<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1275</div><!-- fragment -->
1276</div>
1277</div>
1278<h2 class="groupheader">Variable Documentation</h2>
1279<a class="anchor" id="a0712735973f172ac9efc7d48a31e47ad"></a>
1280<div class="memitem">
1281<div class="memproto">
1282 <table class="memname">
1283 <tr>
1284 <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1285 </tr>
1286 </table>
1287</div><div class="memdoc">
1288
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1289<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00068">68</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]1290
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1291</div>
1292</div>
1293<a class="anchor" id="a4884a666a1e93fbf8c27bd7d2da3c8bb"></a>
1294<div class="memitem">
1295<div class="memproto">
1296 <table class="memname">
1297 <tr>
1298 <td class="memname">__constant uint4 idx4 = (uint4)(0, 1, 2, 3)</td>
1299 </tr>
1300 </table>
1301</div><div class="memdoc">
1302
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1303<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00069">69</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1304
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1305<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1306
1307</div>
1308</div>
1309<a class="anchor" id="aa1dd94b8d98f1c6d790bdf0fc5de29e9"></a>
1310<div class="memitem">
1311<div class="memproto">
1312 <table class="memname">
1313 <tr>
1314 <td class="memname">__constant uint16 idx__ = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1315 </tr>
1316 </table>
1317</div><div class="memdoc">
1318
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1319<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00063">63</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1320
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1321<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1322
1323</div>
1324</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1325<a class="anchor" id="a538b4b63f40e7b12891774e03a4f0dec"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1326<div class="memitem">
1327<div class="memproto">
1328 <table class="memname">
1329 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1330 <td class="memname">__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( -FLT_MAX )</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1331 </tr>
1332 </table>
1333</div><div class="memdoc">
1334
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1335<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00067">67</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]1336
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1337</div>
1338</div>
1339<a class="anchor" id="af7a08044d0e491a0ee1520a24a107a2b"></a>
1340<div class="memitem">
1341<div class="memproto">
1342 <table class="memname">
1343 <tr>
1344 <td class="memname">__constant DATA_TYPE16 type_min_ = ( DATA_TYPE16 )( -FLT_MAX )</td>
1345 </tr>
1346 </table>
1347</div><div class="memdoc">
1348
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1349<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>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1350
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1351<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00306">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00157">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1352
1353</div>
1354</div>
1355</div><!-- contents -->
1356</div><!-- doc-content -->
1357<!-- start footer part -->
1358<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
1359 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1360 <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>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]1361 <li class="footer">Generated on Wed Aug 29 2018 15:31:38 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1362 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1363 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1364 </ul>
1365</div>
1366</body>
1367</html>