Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 7f09cf7dfba61aad6d6d588b1b096f56bafaad5e / . / documentation / softmax__layer_8cl.xhtml
blob: 727cc6c6910f2bf88c85d9bcfafd4f9e737e7dcc [file] [log] [blame]
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]1<!-- HTML header for doxygen 1.8.15-->
2<!-- Remember to use version doxygen 1.8.15 +-->
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]3<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
4<html xmlns="http://www.w3.org/1999/xhtml">
5<head>
6<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
7<meta http-equiv="X-UA-Compatible" content="IE=9"/>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]8<meta name="generator" content="Doxygen 1.8.15"/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]9<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]10<title>Compute Library: src/core/CL/cl_kernels/softmax_layer.cl File Reference</title>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]11<link href="tabs.css" rel="stylesheet" type="text/css"/>
12<script type="text/javascript" src="jquery.js"></script>
13<script type="text/javascript" src="dynsections.js"></script>
14<link href="navtree.css" rel="stylesheet" type="text/css"/>
15<script type="text/javascript" src="resize.js"></script>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]16<script type="text/javascript" src="navtreedata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]17<script type="text/javascript" src="navtree.js"></script>
18<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]19/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]20 $(document).ready(initResizable);
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]21/* @license-end */</script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]22<link href="search/search.css" rel="stylesheet" type="text/css"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]23<script type="text/javascript" src="search/searchdata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]24<script type="text/javascript" src="search/search.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]25<script type="text/x-mathjax-config">
26 MathJax.Hub.Config({
27 extensions: ["tex2jax.js"],
28 jax: ["input/TeX","output/HTML-CSS"],
29});
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]30</script><script type="text/javascript" async="async" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]31<link href="doxygen.css" rel="stylesheet" type="text/css" />
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]32<link href="stylesheet.css" rel="stylesheet" type="text/css"/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]33</head>
34<body>
35<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
36<div id="titlearea">
37<table cellspacing="0" cellpadding="0">
38 <tbody>
39 <tr style="height: 56px;">
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]40 <img alt="Compute Library" src="https://raw.githubusercontent.com/ARM-software/ComputeLibrary/gh-pages/ACL_logo.png" style="max-width: 100%;margin-top: 15px;margin-left: 10px"/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]41 <td style="padding-left: 0.5em;">
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]42 <div id="projectname">
Jenkins7f09cf72020-01-22 18:08:16 +0000[diff] [blame^]43 &#160;<span id="projectnumber">19.11.1</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 -->
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]51<!-- Generated by Doxygen 1.8.15 -->
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]52<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]53/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]54var searchBox = new SearchBox("searchBox", "search",false,'Search');
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]55/* @license-end */
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]56</script>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]57<script type="text/javascript" src="menudata.js"></script>
58<script type="text/javascript" src="menu.js"></script>
59<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]60/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]61$(function() {
62 initMenu('',true,false,'search.php','Search');
63 $(document).ready(function() { init_search(); });
64});
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]65/* @license-end */</script>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]66<div id="main-nav"></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]67</div><!-- top -->
68<div id="side-nav" class="ui-resizable side-nav-resizable">
69 <div id="nav-tree">
70 <div id="nav-tree-contents">
71 <div id="nav-sync" class="sync"></div>
72 </div>
73 </div>
74 <div id="splitbar" style="-moz-user-select:none;"
75 class="ui-resizable-handle">
76 </div>
77</div>
78<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]79/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]80$(document).ready(function(){initNavTree('softmax__layer_8cl.xhtml','');});
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]81/* @license-end */
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]82</script>
83<div id="doc-content">
84<!-- window showing the filter options -->
85<div id="MSearchSelectWindow"
86 onmouseover="return searchBox.OnSearchSelectShow()"
87 onmouseout="return searchBox.OnSearchSelectHide()"
88 onkeydown="return searchBox.OnSearchSelectKey(event)">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]89</div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]90
91<!-- iframe showing the search results (closed by default) -->
92<div id="MSearchResultsWindow">
93<iframe src="javascript:void(0)" frameborder="0"
94 name="MSearchResults" id="MSearchResults">
95</iframe>
96</div>
97
98<div class="header">
99 <div class="summary">
100<a href="#define-members">Macros</a> &#124;
101<a href="#func-members">Functions</a> &#124;
102<a href="#var-members">Variables</a> </div>
103 <div class="headertitle">
104<div class="title">softmax_layer.cl File Reference</div> </div>
105</div><!--header-->
106<div class="contents">
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]107<div class="textblock"><code>#include &quot;<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]108</div>
109<p><a href="softmax__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
110<table class="memberdecls">
111<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
112Macros</h2></td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]113<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]114<tr class="separator:abaa48ad818c44e415fd3f9dd0f27bf01"><td class="memSeparator" colspan="2">&#160;</td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]115<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]116<tr class="separator:a44206a4e5783c7aabacec88aad878c88"><td class="memSeparator" colspan="2">&#160;</td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]117<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]118<tr class="separator:ac3af2d18008cbbf7247ae48fcd6e0c4e"><td class="memSeparator" colspan="2">&#160;</td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]119<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]120<tr class="separator:a22303c4047ec5027c1538d53964b9d0d"><td class="memSeparator" colspan="2">&#160;</td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]121<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]122<tr class="separator:a8cde99b1ce0f3c1dacd49261b0cf03d8"><td class="memSeparator" colspan="2">&#160;</td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]123<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]124<tr class="separator:a93cf800667317d96574477b9f0a75234"><td class="memSeparator" colspan="2">&#160;</td></tr>
125<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]126<tr class="separator:a80b22c555ddadb47cc6ca338a9c49126"><td class="memSeparator" colspan="2">&#160;</td></tr>
127<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>
128<tr class="separator:af5b2e33e3c5fcaab3a213f26c2300170"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]129<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>
130<tr class="separator:a08246606c233e7785a497c09672f366f"><td class="memSeparator" colspan="2">&#160;</td></tr>
131<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>
132<tr class="separator:a7c78836761fa3b5b124efea237dac70f"><td class="memSeparator" colspan="2">&#160;</td></tr>
133<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>
134<tr class="separator:a372393c380805985b813dbb16d589a64"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]135</table><table class="memberdecls">
136<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
137Functions</h2></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]138<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]139<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]140<tr class="separator:ac4247ac0991e85965b7ded764e78f12c"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]141<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>
142<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>
143<tr class="separator:a9d74601bcabbb2f14bcf52385ad666dc"><td class="memSeparator" colspan="2">&#160;</td></tr>
144<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>
145<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>
146<tr class="separator:a82bc360279fcf7cf8033dca6022f21f5"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]147</table><table class="memberdecls">
148<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="var-members"></a>
149Variables</h2></td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]150<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>
151<tr class="separator:af7a08044d0e491a0ee1520a24a107a2b"><td class="memSeparator" colspan="2">&#160;</td></tr>
152<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>
153<tr class="separator:aa1dd94b8d98f1c6d790bdf0fc5de29e9"><td class="memSeparator" colspan="2">&#160;</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]154<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>
155<tr class="separator:a538b4b63f40e7b12891774e03a4f0dec"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]156<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>
157<tr class="separator:a0712735973f172ac9efc7d48a31e47ad"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]158<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>
159<tr class="separator:a4884a666a1e93fbf8c27bd7d2da3c8bb"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]160</table>
161<h2 class="groupheader">Macro Definition Documentation</h2>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]162<a id="a44206a4e5783c7aabacec88aad878c88"></a>
163<h2 class="memtitle"><span class="permalink"><a href="#a44206a4e5783c7aabacec88aad878c88">&#9670;&nbsp;</a></span>ADD_OP</h2>
164
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]165<div class="memitem">
166<div class="memproto">
167 <table class="memname">
168 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]169 <td class="memname">#define ADD_OP</td>
170 <td>(</td>
171 <td class="paramtype">&#160;</td>
172 <td class="paramname">x, </td>
173 </tr>
174 <tr>
175 <td class="paramkey"></td>
176 <td></td>
177 <td class="paramtype">&#160;</td>
178 <td class="paramname">y, </td>
179 </tr>
180 <tr>
181 <td class="paramkey"></td>
182 <td></td>
183 <td class="paramtype">&#160;</td>
184 <td class="paramname">type, </td>
185 </tr>
186 <tr>
187 <td class="paramkey"></td>
188 <td></td>
189 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]190 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]191 </tr>
192 <tr>
193 <td></td>
194 <td>)</td>
195 <td></td><td>&#160;&#160;&#160;((x) + (y))</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]196 </tr>
197 </table>
198</div><div class="memdoc">
199
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]200<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00027">27</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]201
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]202</div>
203</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]204<a id="a8cde99b1ce0f3c1dacd49261b0cf03d8"></a>
205<h2 class="memtitle"><span class="permalink"><a href="#a8cde99b1ce0f3c1dacd49261b0cf03d8">&#9670;&nbsp;</a></span>DIV_OP</h2>
206
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]207<div class="memitem">
208<div class="memproto">
209 <table class="memname">
210 <tr>
211 <td class="memname">#define DIV_OP</td>
212 <td>(</td>
213 <td class="paramtype">&#160;</td>
214 <td class="paramname">x, </td>
215 </tr>
216 <tr>
217 <td class="paramkey"></td>
218 <td></td>
219 <td class="paramtype">&#160;</td>
220 <td class="paramname">y, </td>
221 </tr>
222 <tr>
223 <td class="paramkey"></td>
224 <td></td>
225 <td class="paramtype">&#160;</td>
226 <td class="paramname">type, </td>
227 </tr>
228 <tr>
229 <td class="paramkey"></td>
230 <td></td>
231 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]232 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]233 </tr>
234 <tr>
235 <td></td>
236 <td>)</td>
237 <td></td><td>&#160;&#160;&#160;((x) / (y))</td>
238 </tr>
239 </table>
240</div><div class="memdoc">
241
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]242<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00030">30</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]243
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]244</div>
245</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]246<a id="a93cf800667317d96574477b9f0a75234"></a>
247<h2 class="memtitle"><span class="permalink"><a href="#a93cf800667317d96574477b9f0a75234">&#9670;&nbsp;</a></span>EXP_OP</h2>
248
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]249<div class="memitem">
250<div class="memproto">
251 <table class="memname">
252 <tr>
253 <td class="memname">#define EXP_OP</td>
254 <td>(</td>
255 <td class="paramtype">&#160;</td>
256 <td class="paramname">x, </td>
257 </tr>
258 <tr>
259 <td class="paramkey"></td>
260 <td></td>
261 <td class="paramtype">&#160;</td>
262 <td class="paramname">type, </td>
263 </tr>
264 <tr>
265 <td class="paramkey"></td>
266 <td></td>
267 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]268 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]269 </tr>
270 <tr>
271 <td></td>
272 <td>)</td>
273 <td></td><td>&#160;&#160;&#160;exp((x))</td>
274 </tr>
275 </table>
276</div><div class="memdoc">
277
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]278<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00031">31</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]279
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]280</div>
281</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]282<a id="a08246606c233e7785a497c09672f366f"></a>
283<h2 class="memtitle"><span class="permalink"><a href="#a08246606c233e7785a497c09672f366f">&#9670;&nbsp;</a></span>GRID_SIZE</h2>
284
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]285<div class="memitem">
286<div class="memproto">
287 <table class="memname">
288 <tr>
289 <td class="memname">#define GRID_SIZE&#160;&#160;&#160;1</td>
290 </tr>
291 </table>
292</div><div class="memdoc">
293
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]294<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">43</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]295
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]296</div>
297</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]298<a id="a372393c380805985b813dbb16d589a64"></a>
299<h2 class="memtitle"><span class="permalink"><a href="#a372393c380805985b813dbb16d589a64">&#9670;&nbsp;</a></span>LOG_VECTOR_SIZE</h2>
300
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]301<div class="memitem">
302<div class="memproto">
303 <table class="memname">
304 <tr>
305 <td class="memname">#define LOG_VECTOR_SIZE&#160;&#160;&#160;4</td>
306 </tr>
307 </table>
308</div><div class="memdoc">
309
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]310<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00061">61</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]311
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]312</div>
313</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]314<a id="abaa48ad818c44e415fd3f9dd0f27bf01"></a>
315<h2 class="memtitle"><span class="permalink"><a href="#abaa48ad818c44e415fd3f9dd0f27bf01">&#9670;&nbsp;</a></span>MAX_OP</h2>
316
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]317<div class="memitem">
318<div class="memproto">
319 <table class="memname">
320 <tr>
321 <td class="memname">#define MAX_OP</td>
322 <td>(</td>
323 <td class="paramtype">&#160;</td>
324 <td class="paramname">x, </td>
325 </tr>
326 <tr>
327 <td class="paramkey"></td>
328 <td></td>
329 <td class="paramtype">&#160;</td>
330 <td class="paramname">y, </td>
331 </tr>
332 <tr>
333 <td class="paramkey"></td>
334 <td></td>
335 <td class="paramtype">&#160;</td>
336 <td class="paramname">type, </td>
337 </tr>
338 <tr>
339 <td class="paramkey"></td>
340 <td></td>
341 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]342 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]343 </tr>
344 <tr>
345 <td></td>
346 <td>)</td>
347 <td></td><td>&#160;&#160;&#160;max((x), (y))</td>
348 </tr>
349 </table>
350</div><div class="memdoc">
351
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]352<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00026">26</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]353
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]354</div>
355</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]356<a id="a80b22c555ddadb47cc6ca338a9c49126"></a>
357<h2 class="memtitle"><span class="permalink"><a href="#a80b22c555ddadb47cc6ca338a9c49126">&#9670;&nbsp;</a></span>MINVAL</h2>
358
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]359<div class="memitem">
360<div class="memproto">
361 <table class="memname">
362 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]363 <td class="memname">#define MINVAL&#160;&#160;&#160;-FLT_MAX</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]364 </tr>
365 </table>
366</div><div class="memdoc">
367
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]368<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00037">37</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]369
370</div>
371</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]372<a id="a22303c4047ec5027c1538d53964b9d0d"></a>
373<h2 class="memtitle"><span class="permalink"><a href="#a22303c4047ec5027c1538d53964b9d0d">&#9670;&nbsp;</a></span>MUL_OP</h2>
374
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]375<div class="memitem">
376<div class="memproto">
377 <table class="memname">
378 <tr>
379 <td class="memname">#define MUL_OP</td>
380 <td>(</td>
381 <td class="paramtype">&#160;</td>
382 <td class="paramname">x, </td>
383 </tr>
384 <tr>
385 <td class="paramkey"></td>
386 <td></td>
387 <td class="paramtype">&#160;</td>
388 <td class="paramname">y, </td>
389 </tr>
390 <tr>
391 <td class="paramkey"></td>
392 <td></td>
393 <td class="paramtype">&#160;</td>
394 <td class="paramname">type, </td>
395 </tr>
396 <tr>
397 <td class="paramkey"></td>
398 <td></td>
399 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]400 <td class="paramname">size&#160;</td>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]401 </tr>
402 <tr>
403 <td></td>
404 <td>)</td>
405 <td></td><td>&#160;&#160;&#160;((x) * (y))</td>
406 </tr>
407 </table>
408</div><div class="memdoc">
409
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]410<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00029">29</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]411
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]412</div>
413</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]414<a id="af5b2e33e3c5fcaab3a213f26c2300170"></a>
415<h2 class="memtitle"><span class="permalink"><a href="#af5b2e33e3c5fcaab3a213f26c2300170">&#9670;&nbsp;</a></span>SELECT_DATA_TYPE</h2>
416
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]417<div class="memitem">
418<div class="memproto">
419 <table class="memname">
420 <tr>
421 <td class="memname">#define SELECT_DATA_TYPE&#160;&#160;&#160;int</td>
422 </tr>
423 </table>
424</div><div class="memdoc">
425
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]426<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00038">38</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]427
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]428</div>
429</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]430<a id="ac3af2d18008cbbf7247ae48fcd6e0c4e"></a>
431<h2 class="memtitle"><span class="permalink"><a href="#ac3af2d18008cbbf7247ae48fcd6e0c4e">&#9670;&nbsp;</a></span>SUB_OP</h2>
432
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]433<div class="memitem">
434<div class="memproto">
435 <table class="memname">
436 <tr>
437 <td class="memname">#define SUB_OP</td>
438 <td>(</td>
439 <td class="paramtype">&#160;</td>
440 <td class="paramname">x, </td>
441 </tr>
442 <tr>
443 <td class="paramkey"></td>
444 <td></td>
445 <td class="paramtype">&#160;</td>
446 <td class="paramname">y, </td>
447 </tr>
448 <tr>
449 <td class="paramkey"></td>
450 <td></td>
451 <td class="paramtype">&#160;</td>
452 <td class="paramname">type, </td>
453 </tr>
454 <tr>
455 <td class="paramkey"></td>
456 <td></td>
457 <td class="paramtype">&#160;</td>
Jenkinsb3a371b2018-05-23 11:36:53 +0100[diff] [blame]458 <td class="paramname">size&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]459 </tr>
460 <tr>
461 <td></td>
462 <td>)</td>
463 <td></td><td>&#160;&#160;&#160;((x) - (y))</td>
464 </tr>
465 </table>
466</div><div class="memdoc">
467
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]468<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00028">28</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]469
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]470</div>
471</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]472<a id="a7c78836761fa3b5b124efea237dac70f"></a>
473<h2 class="memtitle"><span class="permalink"><a href="#a7c78836761fa3b5b124efea237dac70f">&#9670;&nbsp;</a></span>VECTOR_SIZE</h2>
474
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]475<div class="memitem">
476<div class="memproto">
477 <table class="memname">
478 <tr>
479 <td class="memname">#define VECTOR_SIZE&#160;&#160;&#160;16</td>
480 </tr>
481 </table>
482</div><div class="memdoc">
483
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]484<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00060">60</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]485
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]486</div>
487</div>
488<h2 class="groupheader">Function Documentation</h2>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]489<a id="a82bc360279fcf7cf8033dca6022f21f5"></a>
490<h2 class="memtitle"><span class="permalink"><a href="#a82bc360279fcf7cf8033dca6022f21f5">&#9670;&nbsp;</a></span>softmax_layer_max_shift_exp_sum_parallel()</h2>
491
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]492<div class="memitem">
493<div class="memproto">
494 <table class="memname">
495 <tr>
496 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_parallel </td>
497 <td>(</td>
498 <td class="paramtype">__global uchar *&#160;</td>
499 <td class="paramname"><em>src_ptr</em>, </td>
500 </tr>
501 <tr>
502 <td class="paramkey"></td>
503 <td></td>
504 <td class="paramtype">uint&#160;</td>
505 <td class="paramname"><em>src_stride_x</em>, </td>
506 </tr>
507 <tr>
508 <td class="paramkey"></td>
509 <td></td>
510 <td class="paramtype">uint&#160;</td>
511 <td class="paramname"><em>src_step_x</em>, </td>
512 </tr>
513 <tr>
514 <td class="paramkey"></td>
515 <td></td>
516 <td class="paramtype">uint&#160;</td>
517 <td class="paramname"><em>src_stride_y</em>, </td>
518 </tr>
519 <tr>
520 <td class="paramkey"></td>
521 <td></td>
522 <td class="paramtype">uint&#160;</td>
523 <td class="paramname"><em>src_step_y</em>, </td>
524 </tr>
525 <tr>
526 <td class="paramkey"></td>
527 <td></td>
528 <td class="paramtype">uint&#160;</td>
529 <td class="paramname"><em>src_stride_z</em>, </td>
530 </tr>
531 <tr>
532 <td class="paramkey"></td>
533 <td></td>
534 <td class="paramtype">uint&#160;</td>
535 <td class="paramname"><em>src_step_z</em>, </td>
536 </tr>
537 <tr>
538 <td class="paramkey"></td>
539 <td></td>
540 <td class="paramtype">uint&#160;</td>
541 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
542 </tr>
543 <tr>
544 <td class="paramkey"></td>
545 <td></td>
546 <td class="paramtype">__global uchar *&#160;</td>
547 <td class="paramname"><em>maxo_ptr</em>, </td>
548 </tr>
549 <tr>
550 <td class="paramkey"></td>
551 <td></td>
552 <td class="paramtype">uint&#160;</td>
553 <td class="paramname"><em>maxo_stride_x</em>, </td>
554 </tr>
555 <tr>
556 <td class="paramkey"></td>
557 <td></td>
558 <td class="paramtype">uint&#160;</td>
559 <td class="paramname"><em>maxo_step_x</em>, </td>
560 </tr>
561 <tr>
562 <td class="paramkey"></td>
563 <td></td>
564 <td class="paramtype">uint&#160;</td>
565 <td class="paramname"><em>maxo_stride_y</em>, </td>
566 </tr>
567 <tr>
568 <td class="paramkey"></td>
569 <td></td>
570 <td class="paramtype">uint&#160;</td>
571 <td class="paramname"><em>maxo_step_y</em>, </td>
572 </tr>
573 <tr>
574 <td class="paramkey"></td>
575 <td></td>
576 <td class="paramtype">uint&#160;</td>
577 <td class="paramname"><em>maxo_stride_z</em>, </td>
578 </tr>
579 <tr>
580 <td class="paramkey"></td>
581 <td></td>
582 <td class="paramtype">uint&#160;</td>
583 <td class="paramname"><em>maxo_step_z</em>, </td>
584 </tr>
585 <tr>
586 <td class="paramkey"></td>
587 <td></td>
588 <td class="paramtype">uint&#160;</td>
589 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
590 </tr>
591 <tr>
592 <td class="paramkey"></td>
593 <td></td>
594 <td class="paramtype">__global uchar *&#160;</td>
595 <td class="paramname"><em>dst_ptr</em>, </td>
596 </tr>
597 <tr>
598 <td class="paramkey"></td>
599 <td></td>
600 <td class="paramtype">uint&#160;</td>
601 <td class="paramname"><em>dst_stride_x</em>, </td>
602 </tr>
603 <tr>
604 <td class="paramkey"></td>
605 <td></td>
606 <td class="paramtype">uint&#160;</td>
607 <td class="paramname"><em>dst_step_x</em>, </td>
608 </tr>
609 <tr>
610 <td class="paramkey"></td>
611 <td></td>
612 <td class="paramtype">uint&#160;</td>
613 <td class="paramname"><em>dst_stride_y</em>, </td>
614 </tr>
615 <tr>
616 <td class="paramkey"></td>
617 <td></td>
618 <td class="paramtype">uint&#160;</td>
619 <td class="paramname"><em>dst_step_y</em>, </td>
620 </tr>
621 <tr>
622 <td class="paramkey"></td>
623 <td></td>
624 <td class="paramtype">uint&#160;</td>
625 <td class="paramname"><em>dst_stride_z</em>, </td>
626 </tr>
627 <tr>
628 <td class="paramkey"></td>
629 <td></td>
630 <td class="paramtype">uint&#160;</td>
631 <td class="paramname"><em>dst_step_z</em>, </td>
632 </tr>
633 <tr>
634 <td class="paramkey"></td>
635 <td></td>
636 <td class="paramtype">uint&#160;</td>
637 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
638 </tr>
639 <tr>
640 <td class="paramkey"></td>
641 <td></td>
642 <td class="paramtype">__global uchar *&#160;</td>
643 <td class="paramname"><em>sum_ptr</em>, </td>
644 </tr>
645 <tr>
646 <td class="paramkey"></td>
647 <td></td>
648 <td class="paramtype">uint&#160;</td>
649 <td class="paramname"><em>sum_stride_x</em>, </td>
650 </tr>
651 <tr>
652 <td class="paramkey"></td>
653 <td></td>
654 <td class="paramtype">uint&#160;</td>
655 <td class="paramname"><em>sum_step_x</em>, </td>
656 </tr>
657 <tr>
658 <td class="paramkey"></td>
659 <td></td>
660 <td class="paramtype">uint&#160;</td>
661 <td class="paramname"><em>sum_stride_y</em>, </td>
662 </tr>
663 <tr>
664 <td class="paramkey"></td>
665 <td></td>
666 <td class="paramtype">uint&#160;</td>
667 <td class="paramname"><em>sum_step_y</em>, </td>
668 </tr>
669 <tr>
670 <td class="paramkey"></td>
671 <td></td>
672 <td class="paramtype">uint&#160;</td>
673 <td class="paramname"><em>sum_stride_z</em>, </td>
674 </tr>
675 <tr>
676 <td class="paramkey"></td>
677 <td></td>
678 <td class="paramtype">uint&#160;</td>
679 <td class="paramname"><em>sum_step_z</em>, </td>
680 </tr>
681 <tr>
682 <td class="paramkey"></td>
683 <td></td>
684 <td class="paramtype">uint&#160;</td>
685 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
686 </tr>
687 <tr>
688 <td class="paramkey"></td>
689 <td></td>
690 <td class="paramtype">uint&#160;</td>
691 <td class="paramname"><em>width</em>&#160;</td>
692 </tr>
693 <tr>
694 <td></td>
695 <td>)</td>
696 <td></td><td></td>
697 </tr>
698 </table>
699</div><div class="memdoc">
700
701<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>
702<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>
703<dd>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]704In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
705<dd>
706Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
707<dl class="params"><dt>Parameters</dt><dd>
708 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]709 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16/F32 </td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]710 <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>
711 <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>
712 <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>
713 <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>
714 <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>
715 <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>
716 <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>
717 <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>
718 <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>
719 <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>
720 <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>
721 <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>
722 <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>
723 <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>
724 <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>
725 <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>
726 <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>
727 <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>
728 <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>
729 <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>
730 <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>
731 <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>
732 <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>
733 <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>
734 <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>
735 <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>
736 <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>
737 <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>
738 <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>
739 <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>
740 <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>
741 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
742 </table>
743 </dd>
744</dl>
745
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]746<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00325">325</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
747<div class="fragment"><div class="line"><a name="l00331"></a><span class="lineno"> 331</span>&#160;{</div><div class="line"><a name="l00332"></a><span class="lineno"> 332</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00333"></a><span class="lineno"> 333</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00334"></a><span class="lineno"> 334</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> maxo = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(maxo);</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="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>);</div><div class="line"><a name="l00336"></a><span class="lineno"> 336</span>&#160;</div><div class="line"><a name="l00337"></a><span class="lineno"> 337</span>&#160; <span class="keyword">const</span> uint lid = get_local_id(0);</div><div class="line"><a name="l00338"></a><span class="lineno"> 338</span>&#160;</div><div class="line"><a name="l00339"></a><span class="lineno"> 339</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00340"></a><span class="lineno"> 340</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00341"></a><span class="lineno"> 341</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00342"></a><span class="lineno"> 342</span>&#160; beta = (<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4))BETA;</div><div class="line"><a name="l00343"></a><span class="lineno"> 343</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00344"></a><span class="lineno"> 344</span>&#160;</div><div class="line"><a name="l00345"></a><span class="lineno"> 345</span>&#160; <span class="comment">// Define one temporary vector per work-item.</span></div><div class="line"><a name="l00346"></a><span class="lineno"> 346</span>&#160; __local <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4) tmp_local[<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>];</div><div class="line"><a name="l00347"></a><span class="lineno"> 347</span>&#160; __local <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> max_local;</div><div class="line"><a name="l00348"></a><span class="lineno"> 348</span>&#160;</div><div class="line"><a name="l00349"></a><span class="lineno"> 349</span>&#160; __constant <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4) type_min4 = (<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4))(<a class="code" href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a>);</div><div class="line"><a name="l00350"></a><span class="lineno"> 350</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00351"></a><span class="lineno"> 351</span>&#160; max_val_vec = (<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4))type_min4;</div><div class="line"><a name="l00352"></a><span class="lineno"> 352</span>&#160; <span class="comment">// Number of elements per work-item.</span></div><div class="line"><a name="l00353"></a><span class="lineno"> 353</span>&#160; const uint row = width / <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00354"></a><span class="lineno"> 354</span>&#160; <span class="comment">// Number of iterations per work-item.</span></div><div class="line"><a name="l00355"></a><span class="lineno"> 355</span>&#160; const uint width_ = row &gt;&gt; 2;</div><div class="line"><a name="l00356"></a><span class="lineno"> 356</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00357"></a><span class="lineno"> 357</span>&#160; uint i = 0;</div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">for</a>(; i &lt; width_; i++)</div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span>&#160; {</div><div class="line"><a name="l00360"></a><span class="lineno"> 360</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00361"></a><span class="lineno"> 361</span>&#160; data_max = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(4)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00362"></a><span class="lineno"> 362</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data_max, max_val_vec, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00363"></a><span class="lineno"> 363</span>&#160; }</div><div class="line"><a name="l00364"></a><span class="lineno"> 364</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00365"></a><span class="lineno"> 365</span>&#160; <span class="comment">// How many work-items needed to complete the computation.</span></div><div class="line"><a name="l00366"></a><span class="lineno"> 366</span>&#160; <span class="comment">//TODO: Optimize this calculation (avoid %).</span></div><div class="line"><a name="l00367"></a><span class="lineno"> 367</span>&#160; <span class="keywordtype">int</span> boundary_workitems = (width % (<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4)) / 4;</div><div class="line"><a name="l00368"></a><span class="lineno"> 368</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00369"></a><span class="lineno"> 369</span>&#160; {</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00371"></a><span class="lineno"> 371</span>&#160; data_max = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(4)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data_max, max_val_vec, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span>&#160; }</div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>&#160; {</div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>&#160; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>&#160; i--;</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span>&#160; }</div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</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="comment">// Handle non multiple of 4</span></div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span>&#160; data_max = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(4)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, (<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * i * 4) + 4, 0));</div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4)</div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span>&#160; widx = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint4)(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * i * 4) + boundary_workitems * 4 + <a class="code" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4));</div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(<a class="code" href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a>, data_max, widx), <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>&#160; }</div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_VECTOR_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00390"></a><span class="lineno"> 390</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="l00391"></a><span class="lineno"> 391</span>&#160; tmp_local[lid] = max_val_vec;</div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span>&#160;</div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 256)</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>(lid &lt; 128)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 128], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span>&#160; }</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 128)</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>(lid &lt; 64)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 64], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 64)</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>(lid &lt; 32)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 32], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 32)</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>(lid &lt; 16)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 16], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00424"></a><span class="lineno"> 424</span>&#160; }</div><div class="line"><a name="l00425"></a><span class="lineno"> 425</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00426"></a><span class="lineno"> 426</span>&#160; }</div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 16)</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>(lid &lt; 8)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 8], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 8)</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>(lid &lt; 4)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 4], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 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; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 4)</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>(lid &lt; 2)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 2], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="keywordflow">if</span>(lid == 0)</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; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(tmp_local[lid + 1], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00454"></a><span class="lineno"> 454</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, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 2);</div><div class="line"><a name="l00455"></a><span class="lineno"> 455</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, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 1);</div><div class="line"><a name="l00456"></a><span class="lineno"> 456</span>&#160; max_local = max_val_vec.s0;</div><div class="line"><a name="l00457"></a><span class="lineno"> 457</span>&#160; }</div><div class="line"><a name="l00458"></a><span class="lineno"> 458</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</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; <span class="comment">/* Second section */</span></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">// Set sum vector</span></div><div class="line"><a name="l00463"></a><span class="lineno"> 463</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00464"></a><span class="lineno"> 464</span>&#160; sum1D = 0;</div><div class="line"><a name="l00465"></a><span class="lineno"> 465</span>&#160; <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> max_val = max_local;</div><div class="line"><a name="l00466"></a><span class="lineno"> 466</span>&#160;</div><div class="line"><a name="l00467"></a><span class="lineno"> 467</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00468"></a><span class="lineno"> 468</span>&#160; <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">for</a>(i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00469"></a><span class="lineno"> 469</span>&#160; {</div><div class="line"><a name="l00470"></a><span class="lineno"> 470</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00471"></a><span class="lineno"> 471</span>&#160; data = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(4)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00472"></a><span class="lineno"> 472</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00473"></a><span class="lineno"> 473</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00474"></a><span class="lineno"> 474</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, beta, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00475"></a><span class="lineno"> 475</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00476"></a><span class="lineno"> 476</span>&#160;<span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00477"></a><span class="lineno"> 477</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00478"></a><span class="lineno"> 478</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00479"></a><span class="lineno"> 479</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00480"></a><span class="lineno"> 480</span>&#160;<span class="preprocessor">#else </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00481"></a><span class="lineno"> 481</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00482"></a><span class="lineno"> 482</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00483"></a><span class="lineno"> 483</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00484"></a><span class="lineno"> 484</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00485"></a><span class="lineno"> 485</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00488"></a><span class="lineno"> 488</span>&#160; <span class="comment">//TODO: Optimize the calculation (avoid %).</span></div><div class="line"><a name="l00489"></a><span class="lineno"> 489</span>&#160; boundary_workitems = (width % (<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4)) / 4;</div><div class="line"><a name="l00490"></a><span class="lineno"> 490</span>&#160; <span class="keywordflow">if</span>(lid &lt; boundary_workitems)</div><div class="line"><a name="l00491"></a><span class="lineno"> 491</span>&#160; {</div><div class="line"><a name="l00492"></a><span class="lineno"> 492</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00493"></a><span class="lineno"> 493</span>&#160; data = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(4)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00494"></a><span class="lineno"> 494</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00495"></a><span class="lineno"> 495</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00496"></a><span class="lineno"> 496</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, beta, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00497"></a><span class="lineno"> 497</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00498"></a><span class="lineno"> 498</span>&#160;<span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00499"></a><span class="lineno"> 499</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00500"></a><span class="lineno"> 500</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00501"></a><span class="lineno"> 501</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00502"></a><span class="lineno"> 502</span>&#160;<span class="preprocessor">#else </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00503"></a><span class="lineno"> 503</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00504"></a><span class="lineno"> 504</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00505"></a><span class="lineno"> 505</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i * <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * 4, 0));</div><div class="line"><a name="l00506"></a><span class="lineno"> 506</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00507"></a><span class="lineno"> 507</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00508"></a><span class="lineno"> 508</span>&#160; }</div><div class="line"><a name="l00509"></a><span class="lineno"> 509</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00510"></a><span class="lineno"> 510</span>&#160; <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00511"></a><span class="lineno"> 511</span>&#160; {</div><div class="line"><a name="l00512"></a><span class="lineno"> 512</span>&#160; boundary_workitems = <a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>;</div><div class="line"><a name="l00513"></a><span class="lineno"> 513</span>&#160; i--;</div><div class="line"><a name="l00514"></a><span class="lineno"> 514</span>&#160; }</div><div class="line"><a name="l00515"></a><span class="lineno"> 515</span>&#160; <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00516"></a><span class="lineno"> 516</span>&#160; {</div><div class="line"><a name="l00517"></a><span class="lineno"> 517</span>&#160; <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="l00518"></a><span class="lineno"> 518</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4)</div><div class="line"><a name="l00519"></a><span class="lineno"> 519</span>&#160; data = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(4)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, (<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * i * 4) + 4, 0));</div><div class="line"><a name="l00520"></a><span class="lineno"> 520</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00521"></a><span class="lineno"> 521</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00522"></a><span class="lineno"> 522</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, beta, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00523"></a><span class="lineno"> 523</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00524"></a><span class="lineno"> 524</span>&#160;<span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00525"></a><span class="lineno"> 525</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00526"></a><span class="lineno"> 526</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, (<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * i * 4) + 4, 0));</div><div class="line"><a name="l00527"></a><span class="lineno"> 527</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00528"></a><span class="lineno"> 528</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4)</div><div class="line"><a name="l00529"></a><span class="lineno"> 529</span>&#160; widx = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint4)(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * i * 4) + boundary_workitems * 4 + <a class="code" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4));</div><div class="line"><a name="l00530"></a><span class="lineno"> 530</span>&#160; data = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(0, data, widx);</div><div class="line"><a name="l00531"></a><span class="lineno"> 531</span>&#160;<span class="preprocessor">#else </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00532"></a><span class="lineno"> 532</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00533"></a><span class="lineno"> 533</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4)</div><div class="line"><a name="l00534"></a><span class="lineno"> 534</span>&#160; widx = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>(((uint4)(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * i * 4) + boundary_workitems * 4 + <a class="code" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, 4));</div><div class="line"><a name="l00535"></a><span class="lineno"> 535</span>&#160; data = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(0, data, widx);</div><div class="line"><a name="l00536"></a><span class="lineno"> 536</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(4)</div><div class="line"><a name="l00537"></a><span class="lineno"> 537</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, (<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> * i * 4) + 4, 0));</div><div class="line"><a name="l00538"></a><span class="lineno"> 538</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00539"></a><span class="lineno"> 539</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 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;<span class="preprocessor">#endif </span><span class="comment">/* NON_MULTIPLE_OF_VECTOR_SIZE */</span><span class="preprocessor"></span></div><div class="line"><a name="l00542"></a><span class="lineno"> 542</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="l00543"></a><span class="lineno"> 543</span>&#160; tmp_local[lid] = sum1D;</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; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00546"></a><span class="lineno"> 546</span>&#160;</div><div class="line"><a name="l00547"></a><span class="lineno"> 547</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 256)</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; <span class="keywordflow">if</span>(lid &lt; 128)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 128], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00554"></a><span class="lineno"> 554</span>&#160; }</div><div class="line"><a name="l00555"></a><span class="lineno"> 555</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 128)</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; <span class="keywordflow">if</span>(lid &lt; 64)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 64], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00562"></a><span class="lineno"> 562</span>&#160; }</div><div class="line"><a name="l00563"></a><span class="lineno"> 563</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 64)</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; <span class="keywordflow">if</span>(lid &lt; 32)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 32], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00570"></a><span class="lineno"> 570</span>&#160; }</div><div class="line"><a name="l00571"></a><span class="lineno"> 571</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 32)</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; <span class="keywordflow">if</span>(lid &lt; 16)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 16], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 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; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00578"></a><span class="lineno"> 578</span>&#160; }</div><div class="line"><a name="l00579"></a><span class="lineno"> 579</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 16)</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; <span class="keywordflow">if</span>(lid &lt; 8)</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; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 8], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</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; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00586"></a><span class="lineno"> 586</span>&#160; }</div><div class="line"><a name="l00587"></a><span class="lineno"> 587</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 8)</div><div class="line"><a name="l00588"></a><span class="lineno"> 588</span>&#160; {</div><div class="line"><a name="l00589"></a><span class="lineno"> 589</span>&#160; <span class="keywordflow">if</span>(lid &lt; 4)</div><div class="line"><a name="l00590"></a><span class="lineno"> 590</span>&#160; {</div><div class="line"><a name="l00591"></a><span class="lineno"> 591</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 4], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00592"></a><span class="lineno"> 592</span>&#160; }</div><div class="line"><a name="l00593"></a><span class="lineno"> 593</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00594"></a><span class="lineno"> 594</span>&#160; }</div><div class="line"><a name="l00595"></a><span class="lineno"> 595</span>&#160; <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> &gt;= 4)</div><div class="line"><a name="l00596"></a><span class="lineno"> 596</span>&#160; {</div><div class="line"><a name="l00597"></a><span class="lineno"> 597</span>&#160; <span class="keywordflow">if</span>(lid &lt; 2)</div><div class="line"><a name="l00598"></a><span class="lineno"> 598</span>&#160; {</div><div class="line"><a name="l00599"></a><span class="lineno"> 599</span>&#160; tmp_local[lid] = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 2], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00600"></a><span class="lineno"> 600</span>&#160; }</div><div class="line"><a name="l00601"></a><span class="lineno"> 601</span>&#160; barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00602"></a><span class="lineno"> 602</span>&#160; }</div><div class="line"><a name="l00603"></a><span class="lineno"> 603</span>&#160; <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00604"></a><span class="lineno"> 604</span>&#160; {</div><div class="line"><a name="l00605"></a><span class="lineno"> 605</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(tmp_local[lid + 1], tmp_local[lid], <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00606"></a><span class="lineno"> 606</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00607"></a><span class="lineno"> 607</span>&#160; sum1D.s01 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01, sum1D.s23, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 2);</div><div class="line"><a name="l00608"></a><span class="lineno"> 608</span>&#160; sum1D.s0 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0, sum1D.s1, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 1);</div><div class="line"><a name="l00609"></a><span class="lineno"> 609</span>&#160; *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>.ptr) = sum1D.s0;</div><div class="line"><a name="l00610"></a><span class="lineno"> 610</span>&#160; }</div><div class="line"><a name="l00611"></a><span class="lineno"> 611</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00510">helpers.h:510</a></div></div>
748<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00261">helpers.h:261</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]749<div class="ttc" id="softmax__layer_8cl_xhtml_a93cf800667317d96574477b9f0a75234"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a></div><div class="ttdeci">#define EXP_OP(x, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00031">softmax_layer.cl:31</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]750<div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]751<div class="ttc" id="softmax__layer_8cl_xhtml_abaa48ad818c44e415fd3f9dd0f27bf01"><div class="ttname"><a href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a></div><div class="ttdeci">#define MAX_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00026">softmax_layer.cl:26</a></div></div>
Jenkins975dfe12019-09-02 11:47:54 +0100[diff] [blame]752<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#l00065">reduction_operation.cl:65</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]753<div class="ttc" id="softmax__layer_8cl_xhtml_af7a08044d0e491a0ee1520a24a107a2b"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a></div><div class="ttdeci">__constant DATA_TYPE16 type_min_</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00062">softmax_layer.cl:62</a></div></div>
754<div class="ttc" id="softmax__layer_8cl_xhtml_a44206a4e5783c7aabacec88aad878c88"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a></div><div class="ttdeci">#define ADD_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00027">softmax_layer.cl:27</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]755<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a0886942393a3ba0dfefaa7516b159784"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">arm_compute::test::validation::for</a></div><div class="ttdeci">for(size_t k=0;k&lt; _target.size();++k)</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_unstack_8cpp_source.xhtml#l00091">Unstack.cpp:91</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]756<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00323">helpers.h:323</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]757<div class="ttc" id="softmax__layer_8cl_xhtml_af5b2e33e3c5fcaab3a213f26c2300170"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a></div><div class="ttdeci">#define SELECT_DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00038">softmax_layer.cl:38</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]758<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#l00070">softmax_layer.cl:70</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]759<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_absolute_difference_8cpp_source.xhtml#l00102">AbsoluteDifference.cpp:102</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]760<div class="ttc" id="softmax__layer_8cl_xhtml_a22303c4047ec5027c1538d53964b9d0d"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00029">softmax_layer.cl:29</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]761<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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00349">helpers.h:349</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]762<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_af77145fbdc6b0c8931148f5597d9de53"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">arm_compute::test::validation::select</a></div><div class="ttdeci">CLSelect select</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_select_8cpp_source.xhtml#l00164">Select.cpp:164</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]763<div class="ttc" id="softmax__layer_8cl_xhtml_ac3af2d18008cbbf7247ae48fcd6e0c4e"><div class="ttname"><a href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a></div><div class="ttdeci">#define SUB_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00028">softmax_layer.cl:28</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]764<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_acb282042d1edeeaa3cc979a206f78b54"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a></div><div class="ttdeci">#define VSTORE(size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00198">helpers.h:198</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]765<div class="ttc" id="softmax__layer_8cl_xhtml_a80b22c555ddadb47cc6ca338a9c49126"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a></div><div class="ttdeci">#define MINVAL</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00037">softmax_layer.cl:37</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]766<div class="ttc" id="softmax__layer_8cl_xhtml_a08246606c233e7785a497c09672f366f"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a></div><div class="ttdeci">#define GRID_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00043">softmax_layer.cl:43</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]767<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a287e2fc366c312b468382c95bb90f91f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a></div><div class="ttdeci">#define VLOAD(size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00195">helpers.h:195</a></div></div>
Jenkins975dfe12019-09-02 11:47:54 +0100[diff] [blame]768<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]769<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">helpers.h:255</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]770</div><!-- fragment -->
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]771<p class="reference">References <a class="el" href="softmax__layer_8cl_source.xhtml#l00027">ADD_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00261">CONVERT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00323">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="_c_l_2_absolute_difference_8cpp_source.xhtml#l00102">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00031">EXP_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00043">GRID_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00070">idx4</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00026">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00037">MINVAL</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00029">MUL_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00510">offset()</a>, <a class="el" href="_c_l_2_select_8cpp_source.xhtml#l00164">arm_compute::test::validation::select</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00038">SELECT_DATA_TYPE</a>, <a class="el" href="_c_l_2_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00028">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00065">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">type_min_</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">VEC_DATA_TYPE</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00195">VLOAD</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00198">VSTORE</a>.</p>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]772
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]773</div>
774</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]775<a id="a9d74601bcabbb2f14bcf52385ad666dc"></a>
776<h2 class="memtitle"><span class="permalink"><a href="#a9d74601bcabbb2f14bcf52385ad666dc">&#9670;&nbsp;</a></span>softmax_layer_max_shift_exp_sum_serial()</h2>
777
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]778<div class="memitem">
779<div class="memproto">
780 <table class="memname">
781 <tr>
782 <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_serial </td>
783 <td>(</td>
784 <td class="paramtype">__global uchar *&#160;</td>
785 <td class="paramname"><em>src_ptr</em>, </td>
786 </tr>
787 <tr>
788 <td class="paramkey"></td>
789 <td></td>
790 <td class="paramtype">uint&#160;</td>
791 <td class="paramname"><em>src_stride_x</em>, </td>
792 </tr>
793 <tr>
794 <td class="paramkey"></td>
795 <td></td>
796 <td class="paramtype">uint&#160;</td>
797 <td class="paramname"><em>src_step_x</em>, </td>
798 </tr>
799 <tr>
800 <td class="paramkey"></td>
801 <td></td>
802 <td class="paramtype">uint&#160;</td>
803 <td class="paramname"><em>src_stride_y</em>, </td>
804 </tr>
805 <tr>
806 <td class="paramkey"></td>
807 <td></td>
808 <td class="paramtype">uint&#160;</td>
809 <td class="paramname"><em>src_step_y</em>, </td>
810 </tr>
811 <tr>
812 <td class="paramkey"></td>
813 <td></td>
814 <td class="paramtype">uint&#160;</td>
815 <td class="paramname"><em>src_stride_z</em>, </td>
816 </tr>
817 <tr>
818 <td class="paramkey"></td>
819 <td></td>
820 <td class="paramtype">uint&#160;</td>
821 <td class="paramname"><em>src_step_z</em>, </td>
822 </tr>
823 <tr>
824 <td class="paramkey"></td>
825 <td></td>
826 <td class="paramtype">uint&#160;</td>
827 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
828 </tr>
829 <tr>
830 <td class="paramkey"></td>
831 <td></td>
832 <td class="paramtype">__global uchar *&#160;</td>
833 <td class="paramname"><em>maxo_ptr</em>, </td>
834 </tr>
835 <tr>
836 <td class="paramkey"></td>
837 <td></td>
838 <td class="paramtype">uint&#160;</td>
839 <td class="paramname"><em>maxo_stride_x</em>, </td>
840 </tr>
841 <tr>
842 <td class="paramkey"></td>
843 <td></td>
844 <td class="paramtype">uint&#160;</td>
845 <td class="paramname"><em>maxo_step_x</em>, </td>
846 </tr>
847 <tr>
848 <td class="paramkey"></td>
849 <td></td>
850 <td class="paramtype">uint&#160;</td>
851 <td class="paramname"><em>maxo_stride_y</em>, </td>
852 </tr>
853 <tr>
854 <td class="paramkey"></td>
855 <td></td>
856 <td class="paramtype">uint&#160;</td>
857 <td class="paramname"><em>maxo_step_y</em>, </td>
858 </tr>
859 <tr>
860 <td class="paramkey"></td>
861 <td></td>
862 <td class="paramtype">uint&#160;</td>
863 <td class="paramname"><em>maxo_stride_z</em>, </td>
864 </tr>
865 <tr>
866 <td class="paramkey"></td>
867 <td></td>
868 <td class="paramtype">uint&#160;</td>
869 <td class="paramname"><em>maxo_step_z</em>, </td>
870 </tr>
871 <tr>
872 <td class="paramkey"></td>
873 <td></td>
874 <td class="paramtype">uint&#160;</td>
875 <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td>
876 </tr>
877 <tr>
878 <td class="paramkey"></td>
879 <td></td>
880 <td class="paramtype">__global uchar *&#160;</td>
881 <td class="paramname"><em>dst_ptr</em>, </td>
882 </tr>
883 <tr>
884 <td class="paramkey"></td>
885 <td></td>
886 <td class="paramtype">uint&#160;</td>
887 <td class="paramname"><em>dst_stride_x</em>, </td>
888 </tr>
889 <tr>
890 <td class="paramkey"></td>
891 <td></td>
892 <td class="paramtype">uint&#160;</td>
893 <td class="paramname"><em>dst_step_x</em>, </td>
894 </tr>
895 <tr>
896 <td class="paramkey"></td>
897 <td></td>
898 <td class="paramtype">uint&#160;</td>
899 <td class="paramname"><em>dst_stride_y</em>, </td>
900 </tr>
901 <tr>
902 <td class="paramkey"></td>
903 <td></td>
904 <td class="paramtype">uint&#160;</td>
905 <td class="paramname"><em>dst_step_y</em>, </td>
906 </tr>
907 <tr>
908 <td class="paramkey"></td>
909 <td></td>
910 <td class="paramtype">uint&#160;</td>
911 <td class="paramname"><em>dst_stride_z</em>, </td>
912 </tr>
913 <tr>
914 <td class="paramkey"></td>
915 <td></td>
916 <td class="paramtype">uint&#160;</td>
917 <td class="paramname"><em>dst_step_z</em>, </td>
918 </tr>
919 <tr>
920 <td class="paramkey"></td>
921 <td></td>
922 <td class="paramtype">uint&#160;</td>
923 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
924 </tr>
925 <tr>
926 <td class="paramkey"></td>
927 <td></td>
928 <td class="paramtype">__global uchar *&#160;</td>
929 <td class="paramname"><em>sum_ptr</em>, </td>
930 </tr>
931 <tr>
932 <td class="paramkey"></td>
933 <td></td>
934 <td class="paramtype">uint&#160;</td>
935 <td class="paramname"><em>sum_stride_x</em>, </td>
936 </tr>
937 <tr>
938 <td class="paramkey"></td>
939 <td></td>
940 <td class="paramtype">uint&#160;</td>
941 <td class="paramname"><em>sum_step_x</em>, </td>
942 </tr>
943 <tr>
944 <td class="paramkey"></td>
945 <td></td>
946 <td class="paramtype">uint&#160;</td>
947 <td class="paramname"><em>sum_stride_y</em>, </td>
948 </tr>
949 <tr>
950 <td class="paramkey"></td>
951 <td></td>
952 <td class="paramtype">uint&#160;</td>
953 <td class="paramname"><em>sum_step_y</em>, </td>
954 </tr>
955 <tr>
956 <td class="paramkey"></td>
957 <td></td>
958 <td class="paramtype">uint&#160;</td>
959 <td class="paramname"><em>sum_stride_z</em>, </td>
960 </tr>
961 <tr>
962 <td class="paramkey"></td>
963 <td></td>
964 <td class="paramtype">uint&#160;</td>
965 <td class="paramname"><em>sum_step_z</em>, </td>
966 </tr>
967 <tr>
968 <td class="paramkey"></td>
969 <td></td>
970 <td class="paramtype">uint&#160;</td>
971 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
972 </tr>
973 <tr>
974 <td class="paramkey"></td>
975 <td></td>
976 <td class="paramtype">uint&#160;</td>
977 <td class="paramname"><em>width</em>&#160;</td>
978 </tr>
979 <tr>
980 <td></td>
981 <td>)</td>
982 <td></td><td></td>
983 </tr>
984 </table>
985</div><div class="memdoc">
986
987<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>
988<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>
989<dd>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]990In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd>
991<dd>
992Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl>
993<dl class="params"><dt>Parameters</dt><dd>
994 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]995 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16/F32 </td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]996 <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>
997 <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>
998 <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>
999 <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>
1000 <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>
1001 <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>
1002 <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>
1003 <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>
1004 <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>
1005 <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>
1006 <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>
1007 <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>
1008 <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>
1009 <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>
1010 <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>
1011 <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>
1012 <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>
1013 <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>
1014 <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>
1015 <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>
1016 <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>
1017 <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>
1018 <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>
1019 <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>
1020 <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>
1021 <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>
1022 <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>
1023 <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>
1024 <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>
1025 <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>
1026 <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>
1027 <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr>
1028 </table>
1029 </dd>
1030</dl>
1031
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1032<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00162">162</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
1033<div class="fragment"><div class="line"><a name="l00168"></a><span class="lineno"> 168</span>&#160;{</div><div class="line"><a name="l00169"></a><span class="lineno"> 169</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00170"></a><span class="lineno"> 170</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00171"></a><span class="lineno"> 171</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> maxo = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(maxo);</div><div class="line"><a name="l00172"></a><span class="lineno"> 172</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>);</div><div class="line"><a name="l00173"></a><span class="lineno"> 173</span>&#160;</div><div class="line"><a name="l00174"></a><span class="lineno"> 174</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00175"></a><span class="lineno"> 175</span>&#160; <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00176"></a><span class="lineno"> 176</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00177"></a><span class="lineno"> 177</span>&#160; beta = (<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>))BETA;</div><div class="line"><a name="l00178"></a><span class="lineno"> 178</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00179"></a><span class="lineno"> 179</span>&#160;</div><div class="line"><a name="l00180"></a><span class="lineno"> 180</span>&#160; <span class="comment">// Initialize local maximum</span></div><div class="line"><a name="l00181"></a><span class="lineno"> 181</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00182"></a><span class="lineno"> 182</span>&#160; max_val_vec = (<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>))<a class="code" href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a>;</div><div class="line"><a name="l00183"></a><span class="lineno"> 183</span>&#160;</div><div class="line"><a name="l00184"></a><span class="lineno"> 184</span>&#160; <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00185"></a><span class="lineno"> 185</span>&#160; <span class="keyword">const</span> uint width_ = width &gt;&gt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>;</div><div class="line"><a name="l00186"></a><span class="lineno"> 186</span>&#160; <span class="keywordflow">for</span>(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00187"></a><span class="lineno"> 187</span>&#160; {</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00189"></a><span class="lineno"> 189</span>&#160; data_max = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00190"></a><span class="lineno"> 190</span>&#160; max_val_vec = <a class="code" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(data_max, max_val_vec, <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="l00191"></a><span class="lineno"> 191</span>&#160; }</div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>&#160;</div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00195"></a><span class="lineno"> 195</span>&#160; data_max = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, width_ &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00196"></a><span class="lineno"> 196</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>&#160; widx = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a>((<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a>(uint, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)))(width_ &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, <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="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(max_val_vec, <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(<a class="code" href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a>, data_max, widx), <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="l00199"></a><span class="lineno"> 199</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="l00200"></a><span class="lineno"> 200</span>&#160;</div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span>&#160; <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00202"></a><span class="lineno"> 202</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00203"></a><span class="lineno"> 203</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, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 8);</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</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="l00205"></a><span class="lineno"> 205</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></div><div class="line"><a name="l00206"></a><span class="lineno"> 206</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, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00207"></a><span class="lineno"> 207</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="l00208"></a><span class="lineno"> 208</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00209"></a><span class="lineno"> 209</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, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 2);</div><div class="line"><a name="l00210"></a><span class="lineno"> 210</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="l00211"></a><span class="lineno"> 211</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, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 1);</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span>&#160; <span class="comment">// Store result</span></div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>&#160; *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)maxo.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = max_val_vec.s0;</div><div class="line"><a name="l00214"></a><span class="lineno"> 214</span>&#160;</div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>&#160; <span class="comment">/* Second section */</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">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>&#160; <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> max_val = *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;maxo, 0, 0));</div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span>&#160;</div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span>&#160; <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00221"></a><span class="lineno"> 221</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00222"></a><span class="lineno"> 222</span>&#160; sum1D = 0;</div><div class="line"><a name="l00223"></a><span class="lineno"> 223</span>&#160;</div><div class="line"><a name="l00224"></a><span class="lineno"> 224</span>&#160; <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00225"></a><span class="lineno"> 225</span>&#160; <span class="keywordflow">for</span>(uint i = 0; i &lt; width_; i++)</div><div class="line"><a name="l00226"></a><span class="lineno"> 226</span>&#160; {</div><div class="line"><a name="l00227"></a><span class="lineno"> 227</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00228"></a><span class="lineno"> 228</span>&#160; data = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00229"></a><span class="lineno"> 229</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>);</div><div class="line"><a name="l00230"></a><span class="lineno"> 230</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00231"></a><span class="lineno"> 231</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, beta, <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="l00232"></a><span class="lineno"> 232</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00233"></a><span class="lineno"> 233</span>&#160;<span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00234"></a><span class="lineno"> 234</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00235"></a><span class="lineno"> 235</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00236"></a><span class="lineno"> 236</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <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="l00237"></a><span class="lineno"> 237</span>&#160;<span class="preprocessor">#else </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00238"></a><span class="lineno"> 238</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <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="l00239"></a><span class="lineno"> 239</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00240"></a><span class="lineno"> 240</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00242"></a><span class="lineno"> 242</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, <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="l00243"></a><span class="lineno"> 243</span>&#160; }</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span>&#160;</div><div class="line"><a name="l00245"></a><span class="lineno"> 245</span>&#160;<span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00247"></a><span class="lineno"> 247</span>&#160; data = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, width_ &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00248"></a><span class="lineno"> 248</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, max_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>);</div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>&#160;<span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(data, beta, <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="l00251"></a><span class="lineno"> 251</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* BETA */</span><span class="preprocessor"></span></div><div class="line"><a name="l00252"></a><span class="lineno"> 252</span>&#160;<span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00254"></a><span class="lineno"> 254</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, width_ &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00255"></a><span class="lineno"> 255</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <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="l00256"></a><span class="lineno"> 256</span>&#160; widx = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a>((<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a>(uint, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)))(width_ &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>));</div><div class="line"><a name="l00257"></a><span class="lineno"> 257</span>&#160; data = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(0, data, widx);</div><div class="line"><a name="l00258"></a><span class="lineno"> 258</span>&#160;<span class="preprocessor">#else </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00259"></a><span class="lineno"> 259</span>&#160; data = <a class="code" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(data, <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="l00260"></a><span class="lineno"> 260</span>&#160; widx = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a>((<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a>((<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a0f6ad555d097377151003a5119ccee45">CL_VEC_DATA_TYPE</a>(uint, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)))(width_ &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) &lt; width, <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>, <a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>));</div><div class="line"><a name="l00261"></a><span class="lineno"> 261</span>&#160; data = <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">select</a>(0, data, widx);</div><div class="line"><a name="l00262"></a><span class="lineno"> 262</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(<a class="code" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>)</div><div class="line"><a name="l00263"></a><span class="lineno"> 263</span>&#160; (data, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, width_ &lt;&lt; <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>, 0));</div><div class="line"><a name="l00264"></a><span class="lineno"> 264</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00265"></a><span class="lineno"> 265</span>&#160; sum1D = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D, data, <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="l00266"></a><span class="lineno"> 266</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="l00267"></a><span class="lineno"> 267</span>&#160;</div><div class="line"><a name="l00268"></a><span class="lineno"> 268</span>&#160; <span class="comment">// Perform sum reduction</span></div><div class="line"><a name="l00269"></a><span class="lineno"> 269</span>&#160;<span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00270"></a><span class="lineno"> 270</span>&#160; sum1D.s01234567 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01234567, sum1D.s89ABCDEF, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 8);</div><div class="line"><a name="l00271"></a><span class="lineno"> 271</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="l00272"></a><span class="lineno"> 272</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 8</span></div><div class="line"><a name="l00273"></a><span class="lineno"> 273</span>&#160; sum1D.s0123 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0123, sum1D.s4567, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 4);</div><div class="line"><a name="l00274"></a><span class="lineno"> 274</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="l00275"></a><span class="lineno"> 275</span>&#160;<span class="preprocessor">#if VECTOR_SIZE &gt;= 4</span></div><div class="line"><a name="l00276"></a><span class="lineno"> 276</span>&#160; sum1D.s01 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s01, sum1D.s23, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 2);</div><div class="line"><a name="l00277"></a><span class="lineno"> 277</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="l00278"></a><span class="lineno"> 278</span>&#160; sum1D.s0 = <a class="code" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(sum1D.s0, sum1D.s1, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 1);</div><div class="line"><a name="l00279"></a><span class="lineno"> 279</span>&#160;</div><div class="line"><a name="l00280"></a><span class="lineno"> 280</span>&#160; <span class="comment">// Calculate and store result</span></div><div class="line"><a name="l00281"></a><span class="lineno"> 281</span>&#160; *((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>.ptr) = sum1D.s0;</div><div class="line"><a name="l00282"></a><span class="lineno"> 282</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00510">helpers.h:510</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1034<div class="ttc" id="softmax__layer_8cl_xhtml_aa1dd94b8d98f1c6d790bdf0fc5de29e9"><div class="ttname"><a href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a></div><div class="ttdeci">__constant uint16 idx__</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00063">softmax_layer.cl:63</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1035<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aa8d95ba04fc73845abc6045952cae5be"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aa8d95ba04fc73845abc6045952cae5be">CONVERT</a></div><div class="ttdeci">#define CONVERT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00261">helpers.h:261</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1036<div class="ttc" id="softmax__layer_8cl_xhtml_a93cf800667317d96574477b9f0a75234"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a></div><div class="ttdeci">#define EXP_OP(x, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00031">softmax_layer.cl:31</a></div></div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1037<div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1038<div class="ttc" id="softmax__layer_8cl_xhtml_a372393c380805985b813dbb16d589a64"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a></div><div class="ttdeci">#define LOG_VECTOR_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00061">softmax_layer.cl:61</a></div></div>
1039<div class="ttc" id="softmax__layer_8cl_xhtml_abaa48ad818c44e415fd3f9dd0f27bf01"><div class="ttname"><a href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a></div><div class="ttdeci">#define MAX_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00026">softmax_layer.cl:26</a></div></div>
Jenkins975dfe12019-09-02 11:47:54 +0100[diff] [blame]1040<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#l00065">reduction_operation.cl:65</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1041<div class="ttc" id="softmax__layer_8cl_xhtml_af7a08044d0e491a0ee1520a24a107a2b"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a></div><div class="ttdeci">__constant DATA_TYPE16 type_min_</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00062">softmax_layer.cl:62</a></div></div>
1042<div class="ttc" id="softmax__layer_8cl_xhtml_a44206a4e5783c7aabacec88aad878c88"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a></div><div class="ttdeci">#define ADD_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00027">softmax_layer.cl:27</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1043<div class="ttc" id="softmax__layer_8cl_xhtml_a7c78836761fa3b5b124efea237dac70f"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a></div><div class="ttdeci">#define VECTOR_SIZE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00060">softmax_layer.cl:60</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1044<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00323">helpers.h:323</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1045<div class="ttc" id="softmax__layer_8cl_xhtml_af5b2e33e3c5fcaab3a213f26c2300170"><div class="ttname"><a href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a></div><div class="ttdeci">#define SELECT_DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00038">softmax_layer.cl:38</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1046<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_ae4b532a93c757194ec73b6790a3e6b1f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#ae4b532a93c757194ec73b6790a3e6b1f">EXPAND</a></div><div class="ttdeci">#define EXPAND(x)</div><div class="ttdoc">Expand the given vector.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00061">helpers.h:61</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]1047<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_absolute_difference_8cpp_source.xhtml#l00102">AbsoluteDifference.cpp:102</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1048<div class="ttc" id="softmax__layer_8cl_xhtml_a22303c4047ec5027c1538d53964b9d0d"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00029">softmax_layer.cl:29</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1049<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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00349">helpers.h:349</a></div></div>
1050<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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00351">helpers.h:351</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]1051<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_af77145fbdc6b0c8931148f5597d9de53"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#af77145fbdc6b0c8931148f5597d9de53">arm_compute::test::validation::select</a></div><div class="ttdeci">CLSelect select</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_select_8cpp_source.xhtml#l00164">Select.cpp:164</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1052<div class="ttc" id="softmax__layer_8cl_xhtml_ac3af2d18008cbbf7247ae48fcd6e0c4e"><div class="ttname"><a href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a></div><div class="ttdeci">#define SUB_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00028">softmax_layer.cl:28</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1053<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_acb282042d1edeeaa3cc979a206f78b54"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a></div><div class="ttdeci">#define VSTORE(size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00198">helpers.h:198</a></div></div>
1054<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a287e2fc366c312b468382c95bb90f91f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a></div><div class="ttdeci">#define VLOAD(size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00195">helpers.h:195</a></div></div>
1055<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a0f6ad555d097377151003a5119ccee45"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00258">helpers.h:258</a></div></div>
Jenkins975dfe12019-09-02 11:47:54 +0100[diff] [blame]1056<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1057<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">helpers.h:255</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1058</div><!-- fragment -->
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1059<p class="reference">References <a class="el" href="softmax__layer_8cl_source.xhtml#l00027">ADD_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00258">CL_VEC_DATA_TYPE</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00261">CONVERT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00323">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="_c_l_2_absolute_difference_8cpp_source.xhtml#l00102">arm_compute::test::validation::dst</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00031">EXP_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00061">EXPAND</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00063">idx__</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00061">LOG_VECTOR_SIZE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00026">MAX_OP</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00029">MUL_OP</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00510">offset()</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00351">Image::ptr</a>, <a class="el" href="_c_l_2_select_8cpp_source.xhtml#l00164">arm_compute::test::validation::select</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00038">SELECT_DATA_TYPE</a>, <a class="el" href="_c_l_2_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00028">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00065">sum()</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">type_min_</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">VEC_DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00060">VECTOR_SIZE</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00195">VLOAD</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00198">VSTORE</a>.</p>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]1060
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1061</div>
1062</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1063<a id="ac4247ac0991e85965b7ded764e78f12c"></a>
1064<h2 class="memtitle"><span class="permalink"><a href="#ac4247ac0991e85965b7ded764e78f12c">&#9670;&nbsp;</a></span>softmax_layer_norm()</h2>
1065
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1066<div class="memitem">
1067<div class="memproto">
1068 <table class="memname">
1069 <tr>
1070 <td class="memname">__kernel void softmax_layer_norm </td>
1071 <td>(</td>
1072 <td class="paramtype">__global uchar *&#160;</td>
1073 <td class="paramname"><em>src_ptr</em>, </td>
1074 </tr>
1075 <tr>
1076 <td class="paramkey"></td>
1077 <td></td>
1078 <td class="paramtype">uint&#160;</td>
1079 <td class="paramname"><em>src_stride_x</em>, </td>
1080 </tr>
1081 <tr>
1082 <td class="paramkey"></td>
1083 <td></td>
1084 <td class="paramtype">uint&#160;</td>
1085 <td class="paramname"><em>src_step_x</em>, </td>
1086 </tr>
1087 <tr>
1088 <td class="paramkey"></td>
1089 <td></td>
1090 <td class="paramtype">uint&#160;</td>
1091 <td class="paramname"><em>src_stride_y</em>, </td>
1092 </tr>
1093 <tr>
1094 <td class="paramkey"></td>
1095 <td></td>
1096 <td class="paramtype">uint&#160;</td>
1097 <td class="paramname"><em>src_step_y</em>, </td>
1098 </tr>
1099 <tr>
1100 <td class="paramkey"></td>
1101 <td></td>
1102 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1103 <td class="paramname"><em>src_stride_z</em>, </td>
1104 </tr>
1105 <tr>
1106 <td class="paramkey"></td>
1107 <td></td>
1108 <td class="paramtype">uint&#160;</td>
1109 <td class="paramname"><em>src_step_z</em>, </td>
1110 </tr>
1111 <tr>
1112 <td class="paramkey"></td>
1113 <td></td>
1114 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1115 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
1116 </tr>
1117 <tr>
1118 <td class="paramkey"></td>
1119 <td></td>
1120 <td class="paramtype">__global uchar *&#160;</td>
1121 <td class="paramname"><em>sum_ptr</em>, </td>
1122 </tr>
1123 <tr>
1124 <td class="paramkey"></td>
1125 <td></td>
1126 <td class="paramtype">uint&#160;</td>
1127 <td class="paramname"><em>sum_stride_x</em>, </td>
1128 </tr>
1129 <tr>
1130 <td class="paramkey"></td>
1131 <td></td>
1132 <td class="paramtype">uint&#160;</td>
1133 <td class="paramname"><em>sum_step_x</em>, </td>
1134 </tr>
1135 <tr>
1136 <td class="paramkey"></td>
1137 <td></td>
1138 <td class="paramtype">uint&#160;</td>
1139 <td class="paramname"><em>sum_stride_y</em>, </td>
1140 </tr>
1141 <tr>
1142 <td class="paramkey"></td>
1143 <td></td>
1144 <td class="paramtype">uint&#160;</td>
1145 <td class="paramname"><em>sum_step_y</em>, </td>
1146 </tr>
1147 <tr>
1148 <td class="paramkey"></td>
1149 <td></td>
1150 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1151 <td class="paramname"><em>sum_stride_z</em>, </td>
1152 </tr>
1153 <tr>
1154 <td class="paramkey"></td>
1155 <td></td>
1156 <td class="paramtype">uint&#160;</td>
1157 <td class="paramname"><em>sum_step_z</em>, </td>
1158 </tr>
1159 <tr>
1160 <td class="paramkey"></td>
1161 <td></td>
1162 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1163 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
1164 </tr>
1165 <tr>
1166 <td class="paramkey"></td>
1167 <td></td>
1168 <td class="paramtype">__global uchar *&#160;</td>
1169 <td class="paramname"><em>dst_ptr</em>, </td>
1170 </tr>
1171 <tr>
1172 <td class="paramkey"></td>
1173 <td></td>
1174 <td class="paramtype">uint&#160;</td>
1175 <td class="paramname"><em>dst_stride_x</em>, </td>
1176 </tr>
1177 <tr>
1178 <td class="paramkey"></td>
1179 <td></td>
1180 <td class="paramtype">uint&#160;</td>
1181 <td class="paramname"><em>dst_step_x</em>, </td>
1182 </tr>
1183 <tr>
1184 <td class="paramkey"></td>
1185 <td></td>
1186 <td class="paramtype">uint&#160;</td>
1187 <td class="paramname"><em>dst_stride_y</em>, </td>
1188 </tr>
1189 <tr>
1190 <td class="paramkey"></td>
1191 <td></td>
1192 <td class="paramtype">uint&#160;</td>
1193 <td class="paramname"><em>dst_step_y</em>, </td>
1194 </tr>
1195 <tr>
1196 <td class="paramkey"></td>
1197 <td></td>
1198 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1199 <td class="paramname"><em>dst_stride_z</em>, </td>
1200 </tr>
1201 <tr>
1202 <td class="paramkey"></td>
1203 <td></td>
1204 <td class="paramtype">uint&#160;</td>
1205 <td class="paramname"><em>dst_step_z</em>, </td>
1206 </tr>
1207 <tr>
1208 <td class="paramkey"></td>
1209 <td></td>
1210 <td class="paramtype">uint&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1211 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
1212 </tr>
1213 <tr>
1214 <td></td>
1215 <td>)</td>
1216 <td></td><td></td>
1217 </tr>
1218 </table>
1219</div><div class="memdoc">
1220
1221<p>Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel. </p>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1222<dl class="section note"><dt>Note</dt><dd>Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short</dd></dl>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1223<dl class="params"><dt>Parameters</dt><dd>
1224 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1225 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor slice. Supported data types: F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1226 <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>
1227 <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>
1228 <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>
1229 <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]1230 <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>
1231 <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]1232 <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]1233 <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]1234 <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>
1235 <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>
1236 <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>
1237 <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]1238 <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>
1239 <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]1240 <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]1241 <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]1242 <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>
1243 <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>
1244 <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>
1245 <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]1246 <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>
1247 <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]1248 <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>
1249 </table>
1250 </dd>
1251</dl>
1252
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1253<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00101">101</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1254<div class="fragment"><div class="line"><a name="l00105"></a><span class="lineno"> 105</span>&#160;{</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>);</div><div class="line"><a name="l00107"></a><span class="lineno"> 107</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00108"></a><span class="lineno"> 108</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a071aa45af973feac43b14f62e54a6fce">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a>(<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>);</div><div class="line"><a name="l00109"></a><span class="lineno"> 109</span>&#160;</div><div class="line"><a name="l00110"></a><span class="lineno"> 110</span>&#160; <span class="comment">// Load max value of 1D logits vector (row)</span></div><div class="line"><a name="l00111"></a><span class="lineno"> 111</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="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>, 0, get_global_id(1)));</div><div class="line"><a name="l00112"></a><span class="lineno"> 112</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00113"></a><span class="lineno"> 113</span>&#160; data = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, 0, 0));</div><div class="line"><a name="l00114"></a><span class="lineno"> 114</span>&#160;<span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span>&#160; vstore16(<a class="code" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(data, sum_val, <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16), 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, 0, 0));</div><div class="line"><a name="l00116"></a><span class="lineno"> 116</span>&#160;<span class="preprocessor">#else </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00117"></a><span class="lineno"> 117</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 <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, 0, 0));</div><div class="line"><a name="l00118"></a><span class="lineno"> 118</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* LOG_SOFTMAX */</span><span class="preprocessor"></span></div><div class="line"><a name="l00119"></a><span class="lineno"> 119</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00510">helpers.h:510</a></div></div>
1255<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a071aa45af973feac43b14f62e54a6fce"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00320">helpers.h:320</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1256<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>
Jenkins975dfe12019-09-02 11:47:54 +0100[diff] [blame]1257<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#l00065">reduction_operation.cl:65</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame]1258<div class="ttc" id="softmax__layer_8cl_xhtml_a8cde99b1ce0f3c1dacd49261b0cf03d8"><div class="ttname"><a href="softmax__layer_8cl.xhtml#a8cde99b1ce0f3c1dacd49261b0cf03d8">DIV_OP</a></div><div class="ttdeci">#define DIV_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00030">softmax_layer.cl:30</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1259<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00323">helpers.h:323</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]1260<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_absolute_difference_8cpp_source.xhtml#l00102">AbsoluteDifference.cpp:102</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1261<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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00349">helpers.h:349</a></div></div>
1262<div class="ttc" id="softmax__layer_8cl_xhtml_ac3af2d18008cbbf7247ae48fcd6e0c4e"><div class="ttname"><a href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a></div><div class="ttdeci">#define SUB_OP(x, y, type, size)</div><div class="ttdef"><b>Definition:</b> <a href="softmax__layer_8cl_source.xhtml#l00028">softmax_layer.cl:28</a></div></div>
Jenkins975dfe12019-09-02 11:47:54 +0100[diff] [blame]1263<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a989ab3e96426615bb98e04e0235088ca"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1264<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">helpers.h:255</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1265</div><!-- fragment -->
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1266<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00323">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00320">CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00030">DIV_OP</a>, <a class="el" href="_c_l_2_absolute_difference_8cpp_source.xhtml#l00102">arm_compute::test::validation::dst</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00510">offset()</a>, <a class="el" href="_c_l_2_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00028">SUB_OP</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00065">sum()</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00255">VEC_DATA_TYPE</a>.</p>
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]1267
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1268</div>
1269</div>
1270<h2 class="groupheader">Variable Documentation</h2>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1271<a id="a0712735973f172ac9efc7d48a31e47ad"></a>
1272<h2 class="memtitle"><span class="permalink"><a href="#a0712735973f172ac9efc7d48a31e47ad">&#9670;&nbsp;</a></span>idx16</h2>
1273
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1274<div class="memitem">
1275<div class="memproto">
1276 <table class="memname">
1277 <tr>
1278 <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1279 </tr>
1280 </table>
1281</div><div class="memdoc">
1282
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1283<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00069">69</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1284
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1285</div>
1286</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1287<a id="a4884a666a1e93fbf8c27bd7d2da3c8bb"></a>
1288<h2 class="memtitle"><span class="permalink"><a href="#a4884a666a1e93fbf8c27bd7d2da3c8bb">&#9670;&nbsp;</a></span>idx4</h2>
1289
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1290<div class="memitem">
1291<div class="memproto">
1292 <table class="memname">
1293 <tr>
1294 <td class="memname">__constant uint4 idx4 = (uint4)(0, 1, 2, 3)</td>
1295 </tr>
1296 </table>
1297</div><div class="memdoc">
1298
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1299<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00070">70</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]1300
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1301<p class="reference">Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00325">softmax_layer_max_shift_exp_sum_parallel()</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1302
1303</div>
1304</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1305<a id="aa1dd94b8d98f1c6d790bdf0fc5de29e9"></a>
1306<h2 class="memtitle"><span class="permalink"><a href="#aa1dd94b8d98f1c6d790bdf0fc5de29e9">&#9670;&nbsp;</a></span>idx__</h2>
1307
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1308<div class="memitem">
1309<div class="memproto">
1310 <table class="memname">
1311 <tr>
1312 <td class="memname">__constant uint16 idx__ = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td>
1313 </tr>
1314 </table>
1315</div><div class="memdoc">
1316
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1317<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00063">63</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1318
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1319<p class="reference">Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00162">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1320
1321</div>
1322</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1323<a id="a538b4b63f40e7b12891774e03a4f0dec"></a>
1324<h2 class="memtitle"><span class="permalink"><a href="#a538b4b63f40e7b12891774e03a4f0dec">&#9670;&nbsp;</a></span>type_min</h2>
1325
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1326<div class="memitem">
1327<div class="memproto">
1328 <table class="memname">
1329 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]1330 <td class="memname">__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( -FLT_MAX )</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1331 </tr>
1332 </table>
1333</div><div class="memdoc">
1334
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1335<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00068">68</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1336
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1337</div>
1338</div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1339<a id="af7a08044d0e491a0ee1520a24a107a2b"></a>
1340<h2 class="memtitle"><span class="permalink"><a href="#af7a08044d0e491a0ee1520a24a107a2b">&#9670;&nbsp;</a></span>type_min_</h2>
1341
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1342<div class="memitem">
1343<div class="memproto">
1344 <table class="memname">
1345 <tr>
1346 <td class="memname">__constant DATA_TYPE16 type_min_ = ( DATA_TYPE16 )( -FLT_MAX )</td>
1347 </tr>
1348 </table>
1349</div><div class="memdoc">
1350
Jenkinsb9abeae2018-11-22 11:58:08 +0000[diff] [blame]1351<p class="definition">Definition at line <a class="el" href="softmax__layer_8cl_source.xhtml#l00062">62</a> of file <a class="el" href="softmax__layer_8cl_source.xhtml">softmax_layer.cl</a>.</p>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1352
Jenkins0e205f72019-11-28 16:53:35 +0000[diff] [blame]1353<p class="reference">Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00325">softmax_layer_max_shift_exp_sum_parallel()</a>, and <a class="el" href="softmax__layer_8cl_source.xhtml#l00162">softmax_layer_max_shift_exp_sum_serial()</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1354
1355</div>
1356</div>
1357</div><!-- contents -->
1358</div><!-- doc-content -->
1359<!-- start footer part -->
1360<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
1361 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]1362 <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>
Jenkins7f09cf72020-01-22 18:08:16 +0000[diff] [blame^]1363 <li class="footer">Generated on Wed Jan 22 2020 18:07:48 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1364 <a href="http://www.doxygen.org/index.html">
Jenkins514be652019-02-28 12:25:18 +0000[diff] [blame]1365 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.15 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1366 </ul>
1367</div>
1368</body>
1369</html>