Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / b3a371bc429d2ba45e56baaf239d8200c2662a74 / . / documentation / softmax__layer_8cl.xhtml
blob: 49401e04f1f463423950e4345d388af9425a5968 [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
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]43 &#160;<span id="projectnumber">18.05</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
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]215<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">44</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]216
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]217<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]257<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00047">47</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]258
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]259<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00120">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
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]293<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00048">48</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]294
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]295<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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
309<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>
310
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]311<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">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
325<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00080">80</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
326
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]327<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]367<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">43</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]368
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]369<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]383<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00054">54</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
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
423<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00046">46</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
424
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]425<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]439<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">55</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]440
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]441<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]481<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">45</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]482
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]483<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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
497<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00079">79</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
498
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]499<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">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>
717Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
718<dd>
719In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
720<dd>
721Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
722<dl class="params"><dt>Parameters</dt><dd>
723 <table class="params">
724 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: QS8/QS16/F16/F32 </td></tr>
725 <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>
726 <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>
727 <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>
728 <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>
729 <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>
730 <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>
731 <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>
732 <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>
733 <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>
734 <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>
735 <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>
736 <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>
737 <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>
738 <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>
739 <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>
740 <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>
741 <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>
742 <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>
743 <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>
744 <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>
745 <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>
746 <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>
747 <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>
748 <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>
749 <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>
750 <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>
751 <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>
752 <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>
753 <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>
754 <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>
755 <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>
756 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
757 </table>
758 </dd>
759</dl>
760
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]761<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">328</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]762
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]763<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00115">arm_compute::test::validation::beta</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00116">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00137">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00048">EXP_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">GRID_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00088">idx4</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00046">MUL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00303">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00144">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">SELECT_DATA_TYPE</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00147">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00081">type_min_</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00042">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00045">VSTORE</a>.</p>
Jenkinsc3f34a42018-03-02 12:38:09 +0000[diff] [blame]764<div class="fragment"><div class="line"><a name="l00334"></a><span class="lineno"> 334</span>&#160;{</div><div class="line"><a name="l00335"></a><span class="lineno"> 335</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="l00336"></a><span class="lineno"> 336</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="l00337"></a><span class="lineno"> 337</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="l00338"></a><span class="lineno"> 338</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="l00339"></a><span class="lineno"> 339</span>&#160;</div><div class="line"><a name="l00340"></a><span class="lineno"> 340</span>&#160; <span class="keyword">const</span> uint lid = get_local_id(0);</div><div class="line"><a name="l00341"></a><span class="lineno"> 341</span>&#160;</div><div class="line"><a name="l00342"></a><span class="lineno"> 342</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00343"></a><span class="lineno"> 343</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00344"></a><span class="lineno"> 344</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00345"></a><span class="lineno"> 345</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a> = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4))BETA;</div><div class="line"><a name="l00346"></a><span class="lineno"> 346</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00347"></a><span class="lineno"> 347</span>&#160;</div><div class="line"><a name="l00348"></a><span class="lineno"> 348</span>&#160; <span class="comment">// Define one temporary vector per work-item.</span></div><div class="line"><a name="l00349"></a><span class="lineno"> 349</span>&#160; __local <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4) tmp_local[<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>];</div><div class="line"><a name="l00350"></a><span class="lineno"> 350</span>&#160; __local DATA_TYPE max_local;</div><div class="line"><a name="l00351"></a><span class="lineno"> 351</span>&#160;</div><div class="line"><a name="l00352"></a><span class="lineno"> 352</span>&#160; __constant <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4) type_min4 = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4))(<a class="code" href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a>);</div><div class="line"><a name="l00353"></a><span class="lineno"> 353</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00354"></a><span class="lineno"> 354</span>&#160; max_val_vec = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4))type_min4;</div><div class="line"><a name="l00355"></a><span class="lineno"> 355</span>&#160; <span class="comment">// Number of elements per work-item.</span></div><div class="line"><a name="l00356"></a><span class="lineno"> 356</span>&#160; const uint row = width / GRID_SIZE;</div><div class="line"><a name="l00357"></a><span class="lineno"> 357</span>&#160; <span class="comment">// Number of iterations per work-item.</span></div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>&#160; const uint width_ = row &gt;&gt; 2;</div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00360"></a><span class="lineno"> 360</span>&#160; uint i = 0;</div><div class="line"><a name="l00361"></a><span class="lineno"> 361</span>&#160; for(; i &lt; width_; i++)</div><div class="line"><a name="l00362"></a><span class="lineno"> 362</span>&#160; {</div><div class="line"><a name="l00363"></a><span class="lineno"> 363</span>&#160; <a class="code" href="fixed__point_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 = <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="l00365"></a><span class="lineno"> 365</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="l00366"></a><span class="lineno"> 366</span>&#160; }</div><div class="line"><a name="l00367"></a><span class="lineno"> 367</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00368"></a><span class="lineno"> 368</span>&#160; <span class="comment">// How many work-items needed to complete the computation.</span></div><div class="line"><a name="l00369"></a><span class="lineno"> 369</span>&#160; <span class="keywordtype">int</span> boundary_workitems = (width % (GRID_SIZE * 4)) / 4;</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00371"></a><span class="lineno"> 371</span>&#160; {</div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</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="l00374"></a><span class="lineno"> 374</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="l00375"></a><span class="lineno"> 375</span>&#160; }</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</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; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>&#160; i--;</div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>&#160; }</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>&#160; {</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span>&#160; <span class="comment">// Handle non multiple of 4</span></div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00386"></a><span class="lineno"> 386</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="l00387"></a><span class="lineno"> 387</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4)</div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint4)(GRID_SIZE * i * 4) + boundary_workitems * 4 + <a class="code" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4));</div><div class="line"><a name="l00389"></a><span class="lineno"> 389</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="l00390"></a><span class="lineno"> 390</span>&#160; }</div><div class="line"><a name="l00391"></a><span class="lineno"> 391</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="l00392"></a><span class="lineno"> 392</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="l00393"></a><span class="lineno"> 393</span>&#160; tmp_local[lid] = max_val_vec;</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; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 256)</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>(lid &lt; 128)</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; 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="l00402"></a><span class="lineno"> 402</span>&#160; }</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>&#160; }</div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 128)</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>(lid &lt; 64)</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; 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="l00410"></a><span class="lineno"> 410</span>&#160; }</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 64)</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>(lid &lt; 32)</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; 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="l00418"></a><span class="lineno"> 418</span>&#160; }</div><div class="line"><a name="l00419"></a><span class="lineno"> 419</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 32)</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>(lid &lt; 16)</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; 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="l00426"></a><span class="lineno"> 426</span>&#160; }</div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 16)</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 &lt; 8)</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; 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="l00434"></a><span class="lineno"> 434</span>&#160; }</div><div class="line"><a name="l00435"></a><span class="lineno"> 435</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00436"></a><span class="lineno"> 436</span>&#160; }</div><div class="line"><a name="l00437"></a><span class="lineno"> 437</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 8)</div><div class="line"><a name="l00438"></a><span class="lineno"> 438</span>&#160; {</div><div class="line"><a name="l00439"></a><span class="lineno"> 439</span>&#160; <span class="keywordflow">if</span>(lid &lt; 4)</div><div class="line"><a name="l00440"></a><span class="lineno"> 440</span>&#160; {</div><div class="line"><a name="l00441"></a><span class="lineno"> 441</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="l00442"></a><span class="lineno"> 442</span>&#160; }</div><div class="line"><a name="l00443"></a><span class="lineno"> 443</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00444"></a><span class="lineno"> 444</span>&#160; }</div><div class="line"><a name="l00445"></a><span class="lineno"> 445</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 4)</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="keywordflow">if</span>(lid &lt; 2)</div><div class="line"><a name="l00448"></a><span class="lineno"> 448</span>&#160; {</div><div class="line"><a name="l00449"></a><span class="lineno"> 449</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="l00450"></a><span class="lineno"> 450</span>&#160; }</div><div class="line"><a name="l00451"></a><span class="lineno"> 451</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00452"></a><span class="lineno"> 452</span>&#160; }</div><div class="line"><a name="l00453"></a><span class="lineno"> 453</span>&#160; <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00454"></a><span class="lineno"> 454</span>&#160; {</div><div class="line"><a name="l00455"></a><span class="lineno"> 455</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="l00456"></a><span class="lineno"> 456</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="l00457"></a><span class="lineno"> 457</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="l00458"></a><span class="lineno"> 458</span>&#160; max_local = max_val_vec.s0;</div><div class="line"><a name="l00459"></a><span class="lineno"> 459</span>&#160; }</div><div class="line"><a name="l00460"></a><span class="lineno"> 460</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00461"></a><span class="lineno"> 461</span>&#160;</div><div class="line"><a name="l00462"></a><span class="lineno"> 462</span>&#160; <span class="comment">/* Second section */</span></div><div class="line"><a name="l00463"></a><span class="lineno"> 463</span>&#160;</div><div class="line"><a name="l00464"></a><span class="lineno"> 464</span>&#160; <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00465"></a><span class="lineno"> 465</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00466"></a><span class="lineno"> 466</span>&#160; sum1D = 0;</div><div class="line"><a name="l00467"></a><span class="lineno"> 467</span>&#160; DATA_TYPE max_val = max_local;</div><div class="line"><a name="l00468"></a><span class="lineno"> 468</span>&#160;</div><div class="line"><a name="l00469"></a><span class="lineno"> 469</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00470"></a><span class="lineno"> 470</span>&#160; for(i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00471"></a><span class="lineno"> 471</span>&#160; {</div><div class="line"><a name="l00472"></a><span class="lineno"> 472</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00473"></a><span class="lineno"> 473</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="l00474"></a><span class="lineno"> 474</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="l00475"></a><span class="lineno"> 475</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00476"></a><span class="lineno"> 476</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="l00477"></a><span class="lineno"> 477</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00478"></a><span class="lineno"> 478</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="l00479"></a><span class="lineno"> 479</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00480"></a><span class="lineno"> 480</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="l00481"></a><span class="lineno"> 481</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="l00482"></a><span class="lineno"> 482</span>&#160; }</div><div class="line"><a name="l00483"></a><span class="lineno"> 483</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00484"></a><span class="lineno"> 484</span>&#160; boundary_workitems = (width % (GRID_SIZE * 4)) / 4;</div><div class="line"><a name="l00485"></a><span class="lineno"> 485</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00486"></a><span class="lineno"> 486</span>&#160; {</div><div class="line"><a name="l00487"></a><span class="lineno"> 487</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00488"></a><span class="lineno"> 488</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="l00489"></a><span class="lineno"> 489</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="l00490"></a><span class="lineno"> 490</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00491"></a><span class="lineno"> 491</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="l00492"></a><span class="lineno"> 492</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00493"></a><span class="lineno"> 493</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="l00494"></a><span class="lineno"> 494</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00495"></a><span class="lineno"> 495</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="l00496"></a><span class="lineno"> 496</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="l00497"></a><span class="lineno"> 497</span>&#160; }</div><div class="line"><a name="l00498"></a><span class="lineno"> 498</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00499"></a><span class="lineno"> 499</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00500"></a><span class="lineno"> 500</span>&#160; {</div><div class="line"><a name="l00501"></a><span class="lineno"> 501</span>&#160; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00502"></a><span class="lineno"> 502</span>&#160; i--;</div><div class="line"><a name="l00503"></a><span class="lineno"> 503</span>&#160; }</div><div class="line"><a name="l00504"></a><span class="lineno"> 504</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00505"></a><span class="lineno"> 505</span>&#160; {</div><div class="line"><a name="l00506"></a><span class="lineno"> 506</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="l00507"></a><span class="lineno"> 507</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, 4)</div><div class="line"><a name="l00508"></a><span class="lineno"> 508</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="l00509"></a><span class="lineno"> 509</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="l00510"></a><span class="lineno"> 510</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00511"></a><span class="lineno"> 511</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="l00512"></a><span class="lineno"> 512</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00513"></a><span class="lineno"> 513</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="l00514"></a><span class="lineno"> 514</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4)</div><div class="line"><a name="l00515"></a><span class="lineno"> 515</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint4)(GRID_SIZE * i * 4) + boundary_workitems * 4 + <a class="code" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, 4));</div><div class="line"><a name="l00516"></a><span class="lineno"> 516</span>&#160; data = select(0, data, widx);</div><div class="line"><a name="l00517"></a><span class="lineno"> 517</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00518"></a><span class="lineno"> 518</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="l00519"></a><span class="lineno"> 519</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="l00520"></a><span class="lineno"> 520</span>&#160; }</div><div class="line"><a name="l00521"></a><span class="lineno"> 521</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="l00522"></a><span class="lineno"> 522</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="l00523"></a><span class="lineno"> 523</span>&#160; tmp_local[lid] = sum1D;</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; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 256)</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>(lid &lt; 128)</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; 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="l00532"></a><span class="lineno"> 532</span>&#160; }</div><div class="line"><a name="l00533"></a><span class="lineno"> 533</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 128)</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>(lid &lt; 64)</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; 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="l00540"></a><span class="lineno"> 540</span>&#160; }</div><div class="line"><a name="l00541"></a><span class="lineno"> 541</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 64)</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>(lid &lt; 32)</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; 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="l00548"></a><span class="lineno"> 548</span>&#160; }</div><div class="line"><a name="l00549"></a><span class="lineno"> 549</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 32)</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>(lid &lt; 16)</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; 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="l00556"></a><span class="lineno"> 556</span>&#160; }</div><div class="line"><a name="l00557"></a><span class="lineno"> 557</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 16)</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 &lt; 8)</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; 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="l00564"></a><span class="lineno"> 564</span>&#160; }</div><div class="line"><a name="l00565"></a><span class="lineno"> 565</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00566"></a><span class="lineno"> 566</span>&#160; }</div><div class="line"><a name="l00567"></a><span class="lineno"> 567</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 8)</div><div class="line"><a name="l00568"></a><span class="lineno"> 568</span>&#160; {</div><div class="line"><a name="l00569"></a><span class="lineno"> 569</span>&#160; <span class="keywordflow">if</span>(lid &lt; 4)</div><div class="line"><a name="l00570"></a><span class="lineno"> 570</span>&#160; {</div><div class="line"><a name="l00571"></a><span class="lineno"> 571</span>&#160; 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="l00572"></a><span class="lineno"> 572</span>&#160; }</div><div class="line"><a name="l00573"></a><span class="lineno"> 573</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00574"></a><span class="lineno"> 574</span>&#160; }</div><div class="line"><a name="l00575"></a><span class="lineno"> 575</span>&#160; <span class="keywordflow">if</span>(GRID_SIZE &gt;= 4)</div><div class="line"><a name="l00576"></a><span class="lineno"> 576</span>&#160; {</div><div class="line"><a name="l00577"></a><span class="lineno"> 577</span>&#160; <span class="keywordflow">if</span>(lid &lt; 2)</div><div class="line"><a name="l00578"></a><span class="lineno"> 578</span>&#160; {</div><div class="line"><a name="l00579"></a><span class="lineno"> 579</span>&#160; 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="l00580"></a><span class="lineno"> 580</span>&#160; }</div><div class="line"><a name="l00581"></a><span class="lineno"> 581</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00582"></a><span class="lineno"> 582</span>&#160; }</div><div class="line"><a name="l00583"></a><span class="lineno"> 583</span>&#160; <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00584"></a><span class="lineno"> 584</span>&#160; {</div><div class="line"><a name="l00585"></a><span class="lineno"> 585</span>&#160; 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="l00586"></a><span class="lineno"> 586</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00587"></a><span class="lineno"> 587</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="l00588"></a><span class="lineno"> 588</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="l00589"></a><span class="lineno"> 589</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div><div class="line"><a name="l00590"></a><span class="lineno"> 590</span>&#160; }</div><div class="line"><a name="l00591"></a><span class="lineno"> 591</span>&#160;}</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#l00115">GEMM.cpp:115</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]765<div class="ttc" id="fixed__point_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00098">fixed_point.h:98</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]766<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#l00042">helpers.h:42</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]767<div class="ttc" id="softmax__layer_8cl_xhtml_a93cf800667317d96574477b9f0a75234"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a></div><div class="ttdeci">#define EXP_OP(x, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00048">softmax_layer.cl:48</a></div></div>
768<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>
769<div class="ttc" id="softmax__layer_8cl_xhtml_abaa48ad818c44e415fd3f9dd0f27bf01"><div class="ttname"><a href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a></div><div class="ttdeci">#define MAX_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00043">softmax_layer.cl:43</a></div></div>
770<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>
771<div class="ttc" id="softmax__layer_8cl_xhtml_af7a08044d0e491a0ee1520a24a107a2b"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a></div><div class="ttdeci">__constant DATA_TYPE16 type_min_</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00081">softmax_layer.cl:81</a></div></div>
772<div class="ttc" id="softmax__layer_8cl_xhtml_a44206a4e5783c7aabacec88aad878c88"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a></div><div class="ttdeci">#define ADD_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00044">softmax_layer.cl:44</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]773<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#l00303">helpers.h:303</a></div></div>
774<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#l00045">helpers.h:45</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]775<div class="ttc" id="softmax__layer_8cl_xhtml_af5b2e33e3c5fcaab3a213f26c2300170"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a></div><div class="ttdeci">#define SELECT_DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00055">softmax_layer.cl:55</a></div></div>
776<div class="ttc" id="softmax__layer_8cl_xhtml_a4884a666a1e93fbf8c27bd7d2da3c8bb"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a></div><div class="ttdeci">__constant uint4 idx4</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00088">softmax_layer.cl:88</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]777<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00137">Convolution.cpp:137</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]778<div class="ttc" id="softmax__layer_8cl_xhtml_a22303c4047ec5027c1538d53964b9d0d"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00046">softmax_layer.cl:46</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]779<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#l00142">helpers.h:142</a></div></div>
780<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#l00144">helpers.h:144</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]781<div class="ttc" id="softmax__layer_8cl_xhtml_ac3af2d18008cbbf7247ae48fcd6e0c4e"><div class="ttname"><a href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a></div><div class="ttdeci">#define SUB_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00045">softmax_layer.cl:45</a></div></div>
782<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
783<div class="ttc" id="softmax__layer_8cl_xhtml_a80b22c555ddadb47cc6ca338a9c49126"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a></div><div class="ttdeci">#define MINVAL</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00054">softmax_layer.cl:54</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]784<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#l00116">helpers.h:116</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]785<div class="ttc" id="softmax__layer_8cl_xhtml_a08246606c233e7785a497c09672f366f"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a></div><div class="ttdeci">#define GRID_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00062">softmax_layer.cl:62</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]786<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]787</div><!-- fragment -->
788</div>
789</div>
790<a class="anchor" id="a9d74601bcabbb2f14bcf52385ad666dc"></a>
791<div class="memitem">
792<div class="memproto">
793 <table class="memname">
794 <tr>
795 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_serial </td>
796 <td>(</td>
797 <td class="paramtype">__global uchar *&#160;</td>
798 <td class="paramname"><em>src_ptr</em>, </td>
799 </tr>
800 <tr>
801 <td class="paramkey"></td>
802 <td></td>
803 <td class="paramtype">uint&#160;</td>
804 <td class="paramname"><em>src_stride_x</em>, </td>
805 </tr>
806 <tr>
807 <td class="paramkey"></td>
808 <td></td>
809 <td class="paramtype">uint&#160;</td>
810 <td class="paramname"><em>src_step_x</em>, </td>
811 </tr>
812 <tr>
813 <td class="paramkey"></td>
814 <td></td>
815 <td class="paramtype">uint&#160;</td>
816 <td class="paramname"><em>src_stride_y</em>, </td>
817 </tr>
818 <tr>
819 <td class="paramkey"></td>
820 <td></td>
821 <td class="paramtype">uint&#160;</td>
822 <td class="paramname"><em>src_step_y</em>, </td>
823 </tr>
824 <tr>
825 <td class="paramkey"></td>
826 <td></td>
827 <td class="paramtype">uint&#160;</td>
828 <td class="paramname"><em>src_stride_z</em>, </td>
829 </tr>
830 <tr>
831 <td class="paramkey"></td>
832 <td></td>
833 <td class="paramtype">uint&#160;</td>
834 <td class="paramname"><em>src_step_z</em>, </td>
835 </tr>
836 <tr>
837 <td class="paramkey"></td>
838 <td></td>
839 <td class="paramtype">uint&#160;</td>
840 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
841 </tr>
842 <tr>
843 <td class="paramkey"></td>
844 <td></td>
845 <td class="paramtype">__global uchar *&#160;</td>
846 <td class="paramname"><em>maxo_ptr</em>, </td>
847 </tr>
848 <tr>
849 <td class="paramkey"></td>
850 <td></td>
851 <td class="paramtype">uint&#160;</td>
852 <td class="paramname"><em>maxo_stride_x</em>, </td>
853 </tr>
854 <tr>
855 <td class="paramkey"></td>
856 <td></td>
857 <td class="paramtype">uint&#160;</td>
858 <td class="paramname"><em>maxo_step_x</em>, </td>
859 </tr>
860 <tr>
861 <td class="paramkey"></td>
862 <td></td>
863 <td class="paramtype">uint&#160;</td>
864 <td class="paramname"><em>maxo_stride_y</em>, </td>
865 </tr>
866 <tr>
867 <td class="paramkey"></td>
868 <td></td>
869 <td class="paramtype">uint&#160;</td>
870 <td class="paramname"><em>maxo_step_y</em>, </td>
871 </tr>
872 <tr>
873 <td class="paramkey"></td>
874 <td></td>
875 <td class="paramtype">uint&#160;</td>
876 <td class="paramname"><em>maxo_stride_z</em>, </td>
877 </tr>
878 <tr>
879 <td class="paramkey"></td>
880 <td></td>
881 <td class="paramtype">uint&#160;</td>
882 <td class="paramname"><em>maxo_step_z</em>, </td>
883 </tr>
884 <tr>
885 <td class="paramkey"></td>
886 <td></td>
887 <td class="paramtype">uint&#160;</td>
888 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
889 </tr>
890 <tr>
891 <td class="paramkey"></td>
892 <td></td>
893 <td class="paramtype">__global uchar *&#160;</td>
894 <td class="paramname"><em>dst_ptr</em>, </td>
895 </tr>
896 <tr>
897 <td class="paramkey"></td>
898 <td></td>
899 <td class="paramtype">uint&#160;</td>
900 <td class="paramname"><em>dst_stride_x</em>, </td>
901 </tr>
902 <tr>
903 <td class="paramkey"></td>
904 <td></td>
905 <td class="paramtype">uint&#160;</td>
906 <td class="paramname"><em>dst_step_x</em>, </td>
907 </tr>
908 <tr>
909 <td class="paramkey"></td>
910 <td></td>
911 <td class="paramtype">uint&#160;</td>
912 <td class="paramname"><em>dst_stride_y</em>, </td>
913 </tr>
914 <tr>
915 <td class="paramkey"></td>
916 <td></td>
917 <td class="paramtype">uint&#160;</td>
918 <td class="paramname"><em>dst_step_y</em>, </td>
919 </tr>
920 <tr>
921 <td class="paramkey"></td>
922 <td></td>
923 <td class="paramtype">uint&#160;</td>
924 <td class="paramname"><em>dst_stride_z</em>, </td>
925 </tr>
926 <tr>
927 <td class="paramkey"></td>
928 <td></td>
929 <td class="paramtype">uint&#160;</td>
930 <td class="paramname"><em>dst_step_z</em>, </td>
931 </tr>
932 <tr>
933 <td class="paramkey"></td>
934 <td></td>
935 <td class="paramtype">uint&#160;</td>
936 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
937 </tr>
938 <tr>
939 <td class="paramkey"></td>
940 <td></td>
941 <td class="paramtype">__global uchar *&#160;</td>
942 <td class="paramname"><em>sum_ptr</em>, </td>
943 </tr>
944 <tr>
945 <td class="paramkey"></td>
946 <td></td>
947 <td class="paramtype">uint&#160;</td>
948 <td class="paramname"><em>sum_stride_x</em>, </td>
949 </tr>
950 <tr>
951 <td class="paramkey"></td>
952 <td></td>
953 <td class="paramtype">uint&#160;</td>
954 <td class="paramname"><em>sum_step_x</em>, </td>
955 </tr>
956 <tr>
957 <td class="paramkey"></td>
958 <td></td>
959 <td class="paramtype">uint&#160;</td>
960 <td class="paramname"><em>sum_stride_y</em>, </td>
961 </tr>
962 <tr>
963 <td class="paramkey"></td>
964 <td></td>
965 <td class="paramtype">uint&#160;</td>
966 <td class="paramname"><em>sum_step_y</em>, </td>
967 </tr>
968 <tr>
969 <td class="paramkey"></td>
970 <td></td>
971 <td class="paramtype">uint&#160;</td>
972 <td class="paramname"><em>sum_stride_z</em>, </td>
973 </tr>
974 <tr>
975 <td class="paramkey"></td>
976 <td></td>
977 <td class="paramtype">uint&#160;</td>
978 <td class="paramname"><em>sum_step_z</em>, </td>
979 </tr>
980 <tr>
981 <td class="paramkey"></td>
982 <td></td>
983 <td class="paramtype">uint&#160;</td>
984 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
985 </tr>
986 <tr>
987 <td class="paramkey"></td>
988 <td></td>
989 <td class="paramtype">uint&#160;</td>
990 <td class="paramname"><em>width</em>&#160;</td>
991 </tr>
992 <tr>
993 <td></td>
994 <td>)</td>
995 <td></td><td></td>
996 </tr>
997 </table>
998</div><div class="memdoc">
999
1000<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>
1001<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>
1002<dd>
1003Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 </dd>
1004<dd>
1005In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
1006<dd>
1007Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
1008<dl class="params"><dt>Parameters</dt><dd>
1009 <table class="params">
1010 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: QS8/QS16/F16/F32 </td></tr>
1011 <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>
1012 <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>
1013 <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>
1014 <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>
1015 <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>
1016 <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>
1017 <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>
1018 <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>
1019 <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>
1020 <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>
1021 <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>
1022 <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>
1023 <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>
1024 <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>
1025 <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>
1026 <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>
1027 <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>
1028 <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>
1029 <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>
1030 <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>
1031 <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>
1032 <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>
1033 <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>
1034 <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>
1035 <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>
1036 <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>
1037 <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>
1038 <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>
1039 <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>
1040 <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>
1041 <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>
1042 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1043 </table>
1044 </dd>
1045</dl>
1046
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1047<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">178</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]1048
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1049<p>References <a class="el" href="softmax__layer_8cl_source.xhtml#l00044">ADD_OP</a>, <a class="el" href="validation_2_c_l_2_g_e_m_m_8cpp_source.xhtml#l00115">arm_compute::test::validation::beta</a>, <a class="el" href="helpers_8h_source.xhtml#l00051">CL_VEC_DATA_TYPE</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00116">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00137">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00048">EXP_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00037">EXPAND</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00082">idx__</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00080">LOG_VECTOR_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00046">MUL_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00303">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00144">Image::ptr</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00055">SELECT_DATA_TYPE</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00147">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00045">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00081">type_min_</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00079">VECTOR_SIZE</a>, <a class="el" href="helpers_8h_source.xhtml#l00042">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00045">VSTORE</a>.</p>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1050<div class="fragment"><div class="line"><a name="l00184"></a><span class="lineno"> 184</span>&#160;{</div><div class="line"><a name="l00185"></a><span class="lineno"> 185</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="l00186"></a><span class="lineno"> 186</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="l00187"></a><span class="lineno"> 187</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="l00188"></a><span class="lineno"> 188</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="l00189"></a><span class="lineno"> 189</span>&#160;</div><div class="line"><a name="l00190"></a><span class="lineno"> 190</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00191"></a><span class="lineno"> 191</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ab146b9cbab6e73e7588b240dc709fe01">beta</a> = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>))BETA;</div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</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">// Initialize local maximum</span></div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00198"></a><span class="lineno"> 198</span>&#160; max_val_vec = (<a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>))<a class="code" href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a>;</div><div class="line"><a name="l00199"></a><span class="lineno"> 199</span>&#160;</div><div class="line"><a name="l00200"></a><span class="lineno"> 200</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00201"></a><span class="lineno"> 201</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="l00202"></a><span class="lineno"> 202</span>&#160; for(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>&#160; {</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00205"></a><span class="lineno"> 205</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="l00206"></a><span class="lineno"> 206</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="l00207"></a><span class="lineno"> 207</span>&#160; }</div><div class="line"><a name="l00208"></a><span class="lineno"> 208</span>&#160;</div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00210"></a><span class="lineno"> 210</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00211"></a><span class="lineno"> 211</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="l00212"></a><span class="lineno"> 212</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, VECTOR_SIZE)</div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((<a class="code" href="helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a>((<a class="code" href="helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a>(uint, VECTOR_SIZE)))(width_ &lt;&lt; LOG_VECTOR_SIZE) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, VECTOR_SIZE));</div><div class="line"><a name="l00214"></a><span class="lineno"> 214</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="l00215"></a><span class="lineno"> 215</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="l00216"></a><span class="lineno"> 216</span>&#160;</div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00219"></a><span class="lineno"> 219</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="l00220"></a><span class="lineno"> 220</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="l00221"></a><span class="lineno"> 221</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></div><div class="line"><a name="l00222"></a><span class="lineno"> 222</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="l00223"></a><span class="lineno"> 223</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="l00224"></a><span class="lineno"> 224</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00225"></a><span class="lineno"> 225</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="l00226"></a><span class="lineno"> 226</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="l00227"></a><span class="lineno"> 227</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="l00228"></a><span class="lineno"> 228</span>&#160; <span class="comment">// Store result</span></div><div class="line"><a name="l00229"></a><span class="lineno"> 229</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="l00230"></a><span class="lineno"> 230</span>&#160;</div><div class="line"><a name="l00231"></a><span class="lineno"> 231</span>&#160; <span class="comment">/* Second section */</span></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; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00234"></a><span class="lineno"> 234</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="l00235"></a><span class="lineno"> 235</span>&#160;</div><div class="line"><a name="l00236"></a><span class="lineno"> 236</span>&#160; <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00237"></a><span class="lineno"> 237</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00238"></a><span class="lineno"> 238</span>&#160; sum1D = 0;</div><div class="line"><a name="l00239"></a><span class="lineno"> 239</span>&#160;</div><div class="line"><a name="l00240"></a><span class="lineno"> 240</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>&#160; for(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00242"></a><span class="lineno"> 242</span>&#160; {</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</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="l00245"></a><span class="lineno"> 245</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="l00246"></a><span class="lineno"> 246</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00247"></a><span class="lineno"> 247</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="l00248"></a><span class="lineno"> 248</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00249"></a><span class="lineno"> 249</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="l00250"></a><span class="lineno"> 250</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VECTOR_SIZE)</div><div class="line"><a name="l00251"></a><span class="lineno"> 251</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="l00252"></a><span class="lineno"> 252</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="l00253"></a><span class="lineno"> 253</span>&#160; }</div><div class="line"><a name="l00254"></a><span class="lineno"> 254</span>&#160;</div><div class="line"><a name="l00255"></a><span class="lineno"> 255</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00256"></a><span class="lineno"> 256</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VECTOR_SIZE)</div><div class="line"><a name="l00257"></a><span class="lineno"> 257</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="l00258"></a><span class="lineno"> 258</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="l00259"></a><span class="lineno"> 259</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00260"></a><span class="lineno"> 260</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="l00261"></a><span class="lineno"> 261</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00262"></a><span class="lineno"> 262</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="l00263"></a><span class="lineno"> 263</span>&#160; widx = <a class="code" href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((<a class="code" href="helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a>((<a class="code" href="helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a>(uint, VECTOR_SIZE)))(width_ &lt;&lt; LOG_VECTOR_SIZE) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) &lt; width, <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(SELECT_DATA_TYPE, VECTOR_SIZE));</div><div class="line"><a name="l00264"></a><span class="lineno"> 264</span>&#160; data = select(0, data, widx);</div><div class="line"><a name="l00265"></a><span class="lineno"> 265</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VECTOR_SIZE)</div><div class="line"><a name="l00266"></a><span class="lineno"> 266</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="l00267"></a><span class="lineno"> 267</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="l00268"></a><span class="lineno"> 268</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="l00269"></a><span class="lineno"> 269</span>&#160;</div><div class="line"><a name="l00270"></a><span class="lineno"> 270</span>&#160; <span class="comment">// Perform sum reduction</span></div><div class="line"><a name="l00271"></a><span class="lineno"> 271</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00272"></a><span class="lineno"> 272</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="l00273"></a><span class="lineno"> 273</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="l00274"></a><span class="lineno"> 274</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></div><div class="line"><a name="l00275"></a><span class="lineno"> 275</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="l00276"></a><span class="lineno"> 276</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="l00277"></a><span class="lineno"> 277</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00278"></a><span class="lineno"> 278</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="l00279"></a><span class="lineno"> 279</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="l00280"></a><span class="lineno"> 280</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="l00281"></a><span class="lineno"> 281</span>&#160;</div><div class="line"><a name="l00282"></a><span class="lineno"> 282</span>&#160; <span class="comment">// Calculate and store result</span></div><div class="line"><a name="l00283"></a><span class="lineno"> 283</span>&#160; *((__global DATA_TYPE *)sum.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = sum1D.s0;</div><div class="line"><a name="l00284"></a><span class="lineno"> 284</span>&#160;}</div><div class="ttc" id="softmax__layer_8cl_xhtml_aa1dd94b8d98f1c6d790bdf0fc5de29e9"><div class="ttname"><a href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a></div><div class="ttdeci">__constant uint16 idx__</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00082">softmax_layer.cl:82</a></div></div>
Jenkinsc3f34a42018-03-02 12:38:09 +0000[diff] [blame]1051<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#l00115">GEMM.cpp:115</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1052<div class="ttc" id="fixed__point_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="fixed__point_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00098">fixed_point.h:98</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1053<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#l00042">helpers.h:42</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1054<div class="ttc" id="softmax__layer_8cl_xhtml_a93cf800667317d96574477b9f0a75234"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a></div><div class="ttdeci">#define EXP_OP(x, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00048">softmax_layer.cl:48</a></div></div>
1055<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>
1056<div class="ttc" id="softmax__layer_8cl_xhtml_a372393c380805985b813dbb16d589a64"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a></div><div class="ttdeci">#define LOG_VECTOR_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00080">softmax_layer.cl:80</a></div></div>
1057<div class="ttc" id="softmax__layer_8cl_xhtml_abaa48ad818c44e415fd3f9dd0f27bf01"><div class="ttname"><a href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a></div><div class="ttdeci">#define MAX_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00043">softmax_layer.cl:43</a></div></div>
1058<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>
1059<div class="ttc" id="softmax__layer_8cl_xhtml_af7a08044d0e491a0ee1520a24a107a2b"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a></div><div class="ttdeci">__constant DATA_TYPE16 type_min_</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00081">softmax_layer.cl:81</a></div></div>
1060<div class="ttc" id="softmax__layer_8cl_xhtml_a44206a4e5783c7aabacec88aad878c88"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a></div><div class="ttdeci">#define ADD_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00044">softmax_layer.cl:44</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1061<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#l00303">helpers.h:303</a></div></div>
1062<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#l00045">helpers.h:45</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1063<div class="ttc" id="softmax__layer_8cl_xhtml_a7c78836761fa3b5b124efea237dac70f"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a></div><div class="ttdeci">#define VECTOR_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00079">softmax_layer.cl:79</a></div></div>
1064<div class="ttc" id="softmax__layer_8cl_xhtml_af5b2e33e3c5fcaab3a213f26c2300170"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a></div><div class="ttdeci">#define SELECT_DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00055">softmax_layer.cl:55</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1065<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00137">Convolution.cpp:137</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1066<div class="ttc" id="softmax__layer_8cl_xhtml_a22303c4047ec5027c1538d53964b9d0d"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00046">softmax_layer.cl:46</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1067<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#l00142">helpers.h:142</a></div></div>
1068<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#l00144">helpers.h:144</a></div></div>
1069<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#l00051">helpers.h:51</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1070<div class="ttc" id="softmax__layer_8cl_xhtml_ac3af2d18008cbbf7247ae48fcd6e0c4e"><div class="ttname"><a href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a></div><div class="ttdeci">#define SUB_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00045">softmax_layer.cl:45</a></div></div>
1071<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1072<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#l00116">helpers.h:116</a></div></div>
1073<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#l00037">helpers.h:37</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1074<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1075</div><!-- fragment -->
1076</div>
1077</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1078<a class="anchor" id="ac4247ac0991e85965b7ded764e78f12c"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1079<div class="memitem">
1080<div class="memproto">
1081 <table class="memname">
1082 <tr>
1083 <td class="memname">__kernel void softmax_layer_norm </td>
1084 <td>(</td>
1085 <td class="paramtype">__global uchar *&#160;</td>
1086 <td class="paramname"><em>src_ptr</em>, </td>
1087 </tr>
1088 <tr>
1089 <td class="paramkey"></td>
1090 <td></td>
1091 <td class="paramtype">uint&#160;</td>
1092 <td class="paramname"><em>src_stride_x</em>, </td>
1093 </tr>
1094 <tr>
1095 <td class="paramkey"></td>
1096 <td></td>
1097 <td class="paramtype">uint&#160;</td>
1098 <td class="paramname"><em>src_step_x</em>, </td>
1099 </tr>
1100 <tr>
1101 <td class="paramkey"></td>
1102 <td></td>
1103 <td class="paramtype">uint&#160;</td>
1104 <td class="paramname"><em>src_stride_y</em>, </td>
1105 </tr>
1106 <tr>
1107 <td class="paramkey"></td>
1108 <td></td>
1109 <td class="paramtype">uint&#160;</td>
1110 <td class="paramname"><em>src_step_y</em>, </td>
1111 </tr>
1112 <tr>
1113 <td class="paramkey"></td>
1114 <td></td>
1115 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1116 <td class="paramname"><em>src_stride_z</em>, </td>
1117 </tr>
1118 <tr>
1119 <td class="paramkey"></td>
1120 <td></td>
1121 <td class="paramtype">uint&#160;</td>
1122 <td class="paramname"><em>src_step_z</em>, </td>
1123 </tr>
1124 <tr>
1125 <td class="paramkey"></td>
1126 <td></td>
1127 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1128 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1129 </tr>
1130 <tr>
1131 <td class="paramkey"></td>
1132 <td></td>
1133 <td class="paramtype">__global uchar *&#160;</td>
1134 <td class="paramname"><em>sum_ptr</em>, </td>
1135 </tr>
1136 <tr>
1137 <td class="paramkey"></td>
1138 <td></td>
1139 <td class="paramtype">uint&#160;</td>
1140 <td class="paramname"><em>sum_stride_x</em>, </td>
1141 </tr>
1142 <tr>
1143 <td class="paramkey"></td>
1144 <td></td>
1145 <td class="paramtype">uint&#160;</td>
1146 <td class="paramname"><em>sum_step_x</em>, </td>
1147 </tr>
1148 <tr>
1149 <td class="paramkey"></td>
1150 <td></td>
1151 <td class="paramtype">uint&#160;</td>
1152 <td class="paramname"><em>sum_stride_y</em>, </td>
1153 </tr>
1154 <tr>
1155 <td class="paramkey"></td>
1156 <td></td>
1157 <td class="paramtype">uint&#160;</td>
1158 <td class="paramname"><em>sum_step_y</em>, </td>
1159 </tr>
1160 <tr>
1161 <td class="paramkey"></td>
1162 <td></td>
1163 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1164 <td class="paramname"><em>sum_stride_z</em>, </td>
1165 </tr>
1166 <tr>
1167 <td class="paramkey"></td>
1168 <td></td>
1169 <td class="paramtype">uint&#160;</td>
1170 <td class="paramname"><em>sum_step_z</em>, </td>
1171 </tr>
1172 <tr>
1173 <td class="paramkey"></td>
1174 <td></td>
1175 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1176 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1177 </tr>
1178 <tr>
1179 <td class="paramkey"></td>
1180 <td></td>
1181 <td class="paramtype">__global uchar *&#160;</td>
1182 <td class="paramname"><em>dst_ptr</em>, </td>
1183 </tr>
1184 <tr>
1185 <td class="paramkey"></td>
1186 <td></td>
1187 <td class="paramtype">uint&#160;</td>
1188 <td class="paramname"><em>dst_stride_x</em>, </td>
1189 </tr>
1190 <tr>
1191 <td class="paramkey"></td>
1192 <td></td>
1193 <td class="paramtype">uint&#160;</td>
1194 <td class="paramname"><em>dst_step_x</em>, </td>
1195 </tr>
1196 <tr>
1197 <td class="paramkey"></td>
1198 <td></td>
1199 <td class="paramtype">uint&#160;</td>
1200 <td class="paramname"><em>dst_stride_y</em>, </td>
1201 </tr>
1202 <tr>
1203 <td class="paramkey"></td>
1204 <td></td>
1205 <td class="paramtype">uint&#160;</td>
1206 <td class="paramname"><em>dst_step_y</em>, </td>
1207 </tr>
1208 <tr>
1209 <td class="paramkey"></td>
1210 <td></td>
1211 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1212 <td class="paramname"><em>dst_stride_z</em>, </td>
1213 </tr>
1214 <tr>
1215 <td class="paramkey"></td>
1216 <td></td>
1217 <td class="paramtype">uint&#160;</td>
1218 <td class="paramname"><em>dst_step_z</em>, </td>
1219 </tr>
1220 <tr>
1221 <td class="paramkey"></td>
1222 <td></td>
1223 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1224 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
1225 </tr>
1226 <tr>
1227 <td></td>
1228 <td>)</td>
1229 <td></td><td></td>
1230 </tr>
1231 </table>
1232</div><div class="memdoc">
1233
1234<p>Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel. </p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1235<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>
1236<dd>
1237Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4</dd></dl>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1238<dl class="params"><dt>Parameters</dt><dd>
1239 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1240 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: QS8/QS16/F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1241 <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>
1242 <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>
1243 <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>
1244 <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]1245 <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>
1246 <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]1247 <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]1248 <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]1249 <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>
1250 <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>
1251 <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>
1252 <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]1253 <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>
1254 <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]1255 <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]1256 <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]1257 <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>
1258 <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>
1259 <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>
1260 <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]1261 <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>
1262 <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]1263 <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>
1264 </table>
1265 </dd>
1266</dl>
1267
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1268<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00120">120</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]1269
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1270<p>References <a class="el" href="helpers_8h_source.xhtml#l00116">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00113">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00047">DIV_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00303">offset()</a>, and <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>.</p>
1271<div class="fragment"><div class="line"><a name="l00124"></a><span class="lineno"> 124</span>&#160;{</div><div class="line"><a name="l00125"></a><span class="lineno"> 125</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="l00126"></a><span class="lineno"> 126</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="l00127"></a><span class="lineno"> 127</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="l00128"></a><span class="lineno"> 128</span>&#160;</div><div class="line"><a name="l00129"></a><span class="lineno"> 129</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00130"></a><span class="lineno"> 130</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="l00131"></a><span class="lineno"> 131</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00132"></a><span class="lineno"> 132</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="l00133"></a><span class="lineno"> 133</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="l00134"></a><span class="lineno"> 134</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#l00113">helpers.h:113</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1272<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>
1273<div class="ttc" id="reduction__operation_8cl_xhtml_ab0df00f5333da51860deb93deb44a782"><div class="ttname"><a href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a></div><div class="ttdeci">DATA_TYPE sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate sum of a vector. </div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00052">reduction_operation.cl:52</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1274<div class="ttc" id="softmax__layer_8cl_xhtml_a8cde99b1ce0f3c1dacd49261b0cf03d8"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a8cde99b1ce0f3c1dacd49261b0cf03d8">DIV_OP</a></div><div class="ttdeci">#define DIV_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00047">softmax_layer.cl:47</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1275<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#l00303">helpers.h:303</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1276<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_adbf67dcee294e673cf796f1ed8aeb6a4"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">arm_compute::test::validation::dst</a></div><div class="ttdeci">CLTensor dst</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00137">Convolution.cpp:137</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1277<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#l00142">helpers.h:142</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1278<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1279<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#l00116">helpers.h:116</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1280<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1281</div><!-- fragment -->
1282</div>
1283</div>
1284<h2 class="groupheader">Variable Documentation</h2>
1285<a class="anchor" id="a0712735973f172ac9efc7d48a31e47ad"></a>
1286<div class="memitem">
1287<div class="memproto">
1288 <table class="memname">
1289 <tr>
1290 <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1291 </tr>
1292 </table>
1293</div><div class="memdoc">
1294
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1295<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00087">87</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1296
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1297</div>
1298</div>
1299<a class="anchor" id="a4884a666a1e93fbf8c27bd7d2da3c8bb"></a>
1300<div class="memitem">
1301<div class="memproto">
1302 <table class="memname">
1303 <tr>
1304 <td class="memname">__constant uint4 idx4 = (uint4)(0, 1, 2, 3)</td>
1305 </tr>
1306 </table>
1307</div><div class="memdoc">
1308
1309<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00088">88</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1310
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1311<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1312
1313</div>
1314</div>
1315<a class="anchor" id="aa1dd94b8d98f1c6d790bdf0fc5de29e9"></a>
1316<div class="memitem">
1317<div class="memproto">
1318 <table class="memname">
1319 <tr>
1320 <td class="memname">__constant uint16 idx__ = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1321 </tr>
1322 </table>
1323</div><div class="memdoc">
1324
1325<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00082">82</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1326
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1327<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1328
1329</div>
1330</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1331<a class="anchor" id="a538b4b63f40e7b12891774e03a4f0dec"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1332<div class="memitem">
1333<div class="memproto">
1334 <table class="memname">
1335 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1336 <td class="memname">__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( -FLT_MAX )</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1337 </tr>
1338 </table>
1339</div><div class="memdoc">
1340
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1341<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00086">86</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1342
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1343</div>
1344</div>
1345<a class="anchor" id="af7a08044d0e491a0ee1520a24a107a2b"></a>
1346<div class="memitem">
1347<div class="memproto">
1348 <table class="memname">
1349 <tr>
1350 <td class="memname">__constant DATA_TYPE16 type_min_ = ( DATA_TYPE16 )( -FLT_MAX )</td>
1351 </tr>
1352 </table>
1353</div><div class="memdoc">
1354
1355<p>Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00081">81</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1356
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame]1357<p>Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00328">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00178">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1358
1359</div>
1360</div>
1361</div><!-- contents -->
1362</div><!-- doc-content -->
1363<!-- start footer part -->
1364<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
1365 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1366 <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>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame^]1367 <li class="footer">Generated on Wed May 23 2018 11:36:39 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1368 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1369 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1370 </ul>
1371</div>
1372</body>
1373</html>