| <!-- HTML header for doxygen 1.8.15--> |
| <!-- Remember to use version doxygen 1.8.15 +--> |
| <!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd"> |
| <html xmlns="http://www.w3.org/1999/xhtml"> |
| <head> |
| <meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/> |
| <meta http-equiv="X-UA-Compatible" content="IE=9"/> |
| <meta name="generator" content="Doxygen 1.8.15"/> |
| <meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines --> |
| <title>Compute Library: src/core/CL/cl_kernels/softmax_layer.cl File Reference</title> |
| <link href="tabs.css" rel="stylesheet" type="text/css"/> |
| <script type="text/javascript" src="jquery.js"></script> |
| <script type="text/javascript" src="dynsections.js"></script> |
| <link href="navtree.css" rel="stylesheet" type="text/css"/> |
| <script type="text/javascript" src="resize.js"></script> |
| <script type="text/javascript" src="navtreedata.js"></script> |
| <script type="text/javascript" src="navtree.js"></script> |
| <script type="text/javascript"> |
| /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&dn=gpl-2.0.txt GPL-v2 */ |
| $(document).ready(initResizable); |
| /* @license-end */</script> |
| <link href="search/search.css" rel="stylesheet" type="text/css"/> |
| <script type="text/javascript" src="search/searchdata.js"></script> |
| <script type="text/javascript" src="search/search.js"></script> |
| <script type="text/x-mathjax-config"> |
| MathJax.Hub.Config({ |
| extensions: ["tex2jax.js"], |
| jax: ["input/TeX","output/HTML-CSS"], |
| }); |
| </script><script type="text/javascript" async="async" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script> |
| <link href="doxygen.css" rel="stylesheet" type="text/css" /> |
| <link href="stylesheet.css" rel="stylesheet" type="text/css"/> |
| </head> |
| <body> |
| <div id="top"><!-- do not remove this div, it is closed by doxygen! --> |
| <div id="titlearea"> |
| <table cellspacing="0" cellpadding="0"> |
| <tbody> |
| <tr style="height: 56px;"> |
| <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"/> |
| <td style="padding-left: 0.5em;"> |
| <div id="projectname"> |
|  <span id="projectnumber">20.02.1</span> |
| </div> |
| </td> |
| </tr> |
| </tbody> |
| </table> |
| </div> |
| <!-- end header part --> |
| <!-- Generated by Doxygen 1.8.15 --> |
| <script type="text/javascript"> |
| /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&dn=gpl-2.0.txt GPL-v2 */ |
| var searchBox = new SearchBox("searchBox", "search",false,'Search'); |
| /* @license-end */ |
| </script> |
| <script type="text/javascript" src="menudata.js"></script> |
| <script type="text/javascript" src="menu.js"></script> |
| <script type="text/javascript"> |
| /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&dn=gpl-2.0.txt GPL-v2 */ |
| $(function() { |
| initMenu('',true,false,'search.php','Search'); |
| $(document).ready(function() { init_search(); }); |
| }); |
| /* @license-end */</script> |
| <div id="main-nav"></div> |
| </div><!-- top --> |
| <div id="side-nav" class="ui-resizable side-nav-resizable"> |
| <div id="nav-tree"> |
| <div id="nav-tree-contents"> |
| <div id="nav-sync" class="sync"></div> |
| </div> |
| </div> |
| <div id="splitbar" style="-moz-user-select:none;" |
| class="ui-resizable-handle"> |
| </div> |
| </div> |
| <script type="text/javascript"> |
| /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&dn=gpl-2.0.txt GPL-v2 */ |
| $(document).ready(function(){initNavTree('softmax__layer_8cl.xhtml','');}); |
| /* @license-end */ |
| </script> |
| <div id="doc-content"> |
| <!-- window showing the filter options --> |
| <div id="MSearchSelectWindow" |
| onmouseover="return searchBox.OnSearchSelectShow()" |
| onmouseout="return searchBox.OnSearchSelectHide()" |
| onkeydown="return searchBox.OnSearchSelectKey(event)"> |
| </div> |
| |
| <!-- iframe showing the search results (closed by default) --> |
| <div id="MSearchResultsWindow"> |
| <iframe src="javascript:void(0)" frameborder="0" |
| name="MSearchResults" id="MSearchResults"> |
| </iframe> |
| </div> |
| |
| <div class="header"> |
| <div class="summary"> |
| <a href="#define-members">Macros</a> | |
| <a href="#func-members">Functions</a> | |
| <a href="#var-members">Variables</a> </div> |
| <div class="headertitle"> |
| <div class="title">softmax_layer.cl File Reference</div> </div> |
| </div><!--header--> |
| <div class="contents"> |
| <div class="textblock"><code>#include "<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml">helpers.h</a>"</code><br /> |
| </div> |
| <p><a href="softmax__layer_8cl_source.xhtml">Go to the source code of this file.</a></p> |
| <table class="memberdecls"> |
| <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a> |
| Macros</h2></td></tr> |
| <tr class="memitem:abaa48ad818c44e415fd3f9dd0f27bf01"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#abaa48ad818c44e415fd3f9dd0f27bf01">MAX_OP</a>(x, y, type, size)   max((x), (y))</td></tr> |
| <tr class="separator:abaa48ad818c44e415fd3f9dd0f27bf01"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a44206a4e5783c7aabacec88aad878c88"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a44206a4e5783c7aabacec88aad878c88">ADD_OP</a>(x, y, type, size)   ((x) + (y))</td></tr> |
| <tr class="separator:a44206a4e5783c7aabacec88aad878c88"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:ac3af2d18008cbbf7247ae48fcd6e0c4e"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#ac3af2d18008cbbf7247ae48fcd6e0c4e">SUB_OP</a>(x, y, type, size)   ((x) - (y))</td></tr> |
| <tr class="separator:ac3af2d18008cbbf7247ae48fcd6e0c4e"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a22303c4047ec5027c1538d53964b9d0d"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a22303c4047ec5027c1538d53964b9d0d">MUL_OP</a>(x, y, type, size)   ((x) * (y))</td></tr> |
| <tr class="separator:a22303c4047ec5027c1538d53964b9d0d"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a8cde99b1ce0f3c1dacd49261b0cf03d8"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a8cde99b1ce0f3c1dacd49261b0cf03d8">DIV_OP</a>(x, y, type, size)   ((x) / (y))</td></tr> |
| <tr class="separator:a8cde99b1ce0f3c1dacd49261b0cf03d8"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a93cf800667317d96574477b9f0a75234"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a93cf800667317d96574477b9f0a75234">EXP_OP</a>(x, type, size)   exp((x))</td></tr> |
| <tr class="separator:a93cf800667317d96574477b9f0a75234"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a80b22c555ddadb47cc6ca338a9c49126"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a80b22c555ddadb47cc6ca338a9c49126">MINVAL</a>   -FLT_MAX</td></tr> |
| <tr class="separator:a80b22c555ddadb47cc6ca338a9c49126"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:af5b2e33e3c5fcaab3a213f26c2300170"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#af5b2e33e3c5fcaab3a213f26c2300170">SELECT_DATA_TYPE</a>   int</td></tr> |
| <tr class="separator:af5b2e33e3c5fcaab3a213f26c2300170"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a08246606c233e7785a497c09672f366f"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a>   1</td></tr> |
| <tr class="separator:a08246606c233e7785a497c09672f366f"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a7c78836761fa3b5b124efea237dac70f"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a7c78836761fa3b5b124efea237dac70f">VECTOR_SIZE</a>   16</td></tr> |
| <tr class="separator:a7c78836761fa3b5b124efea237dac70f"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a372393c380805985b813dbb16d589a64"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>   4</td></tr> |
| <tr class="separator:a372393c380805985b813dbb16d589a64"><td class="memSeparator" colspan="2"> </td></tr> |
| </table><table class="memberdecls"> |
| <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a> |
| Functions</h2></td></tr> |
| <tr class="memitem:ac4247ac0991e85965b7ded764e78f12c"><td class="memItemLeft" align="right" valign="top">__kernel void </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> |
| <tr class="memdesc:ac4247ac0991e85965b7ded764e78f12c"><td class="mdescLeft"> </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> |
| <tr class="separator:ac4247ac0991e85965b7ded764e78f12c"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a9d74601bcabbb2f14bcf52385ad666dc"><td class="memItemLeft" align="right" valign="top">__kernel void </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> |
| <tr class="memdesc:a9d74601bcabbb2f14bcf52385ad666dc"><td class="mdescLeft"> </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> |
| <tr class="separator:a9d74601bcabbb2f14bcf52385ad666dc"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a82bc360279fcf7cf8033dca6022f21f5"><td class="memItemLeft" align="right" valign="top">__kernel void </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> |
| <tr class="memdesc:a82bc360279fcf7cf8033dca6022f21f5"><td class="mdescLeft"> </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> |
| <tr class="separator:a82bc360279fcf7cf8033dca6022f21f5"><td class="memSeparator" colspan="2"> </td></tr> |
| </table><table class="memberdecls"> |
| <tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="var-members"></a> |
| Variables</h2></td></tr> |
| <tr class="memitem:af7a08044d0e491a0ee1520a24a107a2b"><td class="memItemLeft" align="right" valign="top">__constant DATA_TYPE16 </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#af7a08044d0e491a0ee1520a24a107a2b">type_min_</a> = ( DATA_TYPE16 )( -FLT_MAX )</td></tr> |
| <tr class="separator:af7a08044d0e491a0ee1520a24a107a2b"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:aa1dd94b8d98f1c6d790bdf0fc5de29e9"><td class="memItemLeft" align="right" valign="top">__constant uint16 </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> |
| <tr class="separator:aa1dd94b8d98f1c6d790bdf0fc5de29e9"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a538b4b63f40e7b12891774e03a4f0dec"><td class="memItemLeft" align="right" valign="top">__constant DATA_TYPE16 </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a538b4b63f40e7b12891774e03a4f0dec">type_min</a> = ( DATA_TYPE16 )( -FLT_MAX )</td></tr> |
| <tr class="separator:a538b4b63f40e7b12891774e03a4f0dec"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a0712735973f172ac9efc7d48a31e47ad"><td class="memItemLeft" align="right" valign="top">__constant uint16 </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> |
| <tr class="separator:a0712735973f172ac9efc7d48a31e47ad"><td class="memSeparator" colspan="2"> </td></tr> |
| <tr class="memitem:a4884a666a1e93fbf8c27bd7d2da3c8bb"><td class="memItemLeft" align="right" valign="top">__constant uint4 </td><td class="memItemRight" valign="bottom"><a class="el" href="softmax__layer_8cl.xhtml#a4884a666a1e93fbf8c27bd7d2da3c8bb">idx4</a> = (uint4)(0, 1, 2, 3)</td></tr> |
| <tr class="separator:a4884a666a1e93fbf8c27bd7d2da3c8bb"><td class="memSeparator" colspan="2"> </td></tr> |
| </table> |
| <h2 class="groupheader">Macro Definition Documentation</h2> |
| <a id="a44206a4e5783c7aabacec88aad878c88"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a44206a4e5783c7aabacec88aad878c88">◆ </a></span>ADD_OP</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define ADD_OP</td> |
| <td>(</td> |
| <td class="paramtype"> </td> |
| <td class="paramname">x, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">y, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">type, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">size </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td>   ((x) + (y))</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a8cde99b1ce0f3c1dacd49261b0cf03d8"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a8cde99b1ce0f3c1dacd49261b0cf03d8">◆ </a></span>DIV_OP</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define DIV_OP</td> |
| <td>(</td> |
| <td class="paramtype"> </td> |
| <td class="paramname">x, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">y, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">type, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">size </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td>   ((x) / (y))</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a93cf800667317d96574477b9f0a75234"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a93cf800667317d96574477b9f0a75234">◆ </a></span>EXP_OP</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define EXP_OP</td> |
| <td>(</td> |
| <td class="paramtype"> </td> |
| <td class="paramname">x, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">type, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">size </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td>   exp((x))</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a08246606c233e7785a497c09672f366f"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a08246606c233e7785a497c09672f366f">◆ </a></span>GRID_SIZE</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define GRID_SIZE   1</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a372393c380805985b813dbb16d589a64"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a372393c380805985b813dbb16d589a64">◆ </a></span>LOG_VECTOR_SIZE</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define LOG_VECTOR_SIZE   4</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="abaa48ad818c44e415fd3f9dd0f27bf01"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#abaa48ad818c44e415fd3f9dd0f27bf01">◆ </a></span>MAX_OP</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define MAX_OP</td> |
| <td>(</td> |
| <td class="paramtype"> </td> |
| <td class="paramname">x, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">y, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">type, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">size </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td>   max((x), (y))</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a80b22c555ddadb47cc6ca338a9c49126"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a80b22c555ddadb47cc6ca338a9c49126">◆ </a></span>MINVAL</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define MINVAL   -FLT_MAX</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a22303c4047ec5027c1538d53964b9d0d"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a22303c4047ec5027c1538d53964b9d0d">◆ </a></span>MUL_OP</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define MUL_OP</td> |
| <td>(</td> |
| <td class="paramtype"> </td> |
| <td class="paramname">x, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">y, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">type, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">size </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td>   ((x) * (y))</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="af5b2e33e3c5fcaab3a213f26c2300170"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#af5b2e33e3c5fcaab3a213f26c2300170">◆ </a></span>SELECT_DATA_TYPE</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define SELECT_DATA_TYPE   int</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="ac3af2d18008cbbf7247ae48fcd6e0c4e"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#ac3af2d18008cbbf7247ae48fcd6e0c4e">◆ </a></span>SUB_OP</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define SUB_OP</td> |
| <td>(</td> |
| <td class="paramtype"> </td> |
| <td class="paramname">x, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">y, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">type, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype"> </td> |
| <td class="paramname">size </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td>   ((x) - (y))</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a7c78836761fa3b5b124efea237dac70f"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a7c78836761fa3b5b124efea237dac70f">◆ </a></span>VECTOR_SIZE</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">#define VECTOR_SIZE   16</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <h2 class="groupheader">Function Documentation</h2> |
| <a id="a82bc360279fcf7cf8033dca6022f21f5"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a82bc360279fcf7cf8033dca6022f21f5">◆ </a></span>softmax_layer_max_shift_exp_sum_parallel()</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_parallel </td> |
| <td>(</td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>src_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>maxo_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>dst_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>sum_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>width</em> </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td></td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| <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> |
| <dd> |
| In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd> |
| <dd> |
| Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl> |
| <dl class="params"><dt>Parameters</dt><dd> |
| <table class="params"> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr> |
| </table> |
| </dd> |
| </dl> |
| |
| <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> |
| <div class="fragment"><div class="line"><a name="l00331"></a><span class="lineno"> 331</span> {</div><div class="line"><a name="l00332"></a><span class="lineno"> 332</span>  <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>  <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>  <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>  <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> </div><div class="line"><a name="l00337"></a><span class="lineno"> 337</span>  <span class="keyword">const</span> uint lid = get_local_id(0);</div><div class="line"><a name="l00338"></a><span class="lineno"> 338</span> </div><div class="line"><a name="l00339"></a><span class="lineno"> 339</span> <span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00340"></a><span class="lineno"> 340</span>  <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00341"></a><span class="lineno"> 341</span>  <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>  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> <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> </div><div class="line"><a name="l00345"></a><span class="lineno"> 345</span>  <span class="comment">// Define one temporary vector per work-item.</span></div><div class="line"><a name="l00346"></a><span class="lineno"> 346</span>  __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>  __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> </div><div class="line"><a name="l00349"></a><span class="lineno"> 349</span>  __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>  <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>  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>  <span class="comment">// Number of elements per work-item.</span></div><div class="line"><a name="l00353"></a><span class="lineno"> 353</span>  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>  <span class="comment">// Number of iterations per work-item.</span></div><div class="line"><a name="l00355"></a><span class="lineno"> 355</span>  const uint width_ = row >> 2;</div><div class="line"><a name="l00356"></a><span class="lineno"> 356</span>  <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00357"></a><span class="lineno"> 357</span>  uint i = 0;</div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">for</a>(; i < width_; i++)</div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span>  {</div><div class="line"><a name="l00360"></a><span class="lineno"> 360</span>  <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>  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>(&<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>  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>  }</div><div class="line"><a name="l00364"></a><span class="lineno"> 364</span> <span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00365"></a><span class="lineno"> 365</span>  <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>  <span class="comment">//TODO: Optimize this calculation (avoid %).</span></div><div class="line"><a name="l00367"></a><span class="lineno"> 367</span>  <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>  <span class="keywordflow">if</span>(lid < boundary_workitems)</div><div class="line"><a name="l00369"></a><span class="lineno"> 369</span>  {</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span>  <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>  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>(&<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>  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>  }</div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span> <span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>  <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>  {</div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>  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>  i--;</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span>  }</div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>  <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>  {</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>  <span class="comment">// Handle non multiple of 4</span></div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>  <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>  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>(&<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>  <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>  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>) < 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>  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>  }</div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span> <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> <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>  tmp_local[lid] = max_val_vec;</div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span> </div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span> </div><div class="line"><a name="l00395"></a><span class="lineno"> 395</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 256)</div><div class="line"><a name="l00396"></a><span class="lineno"> 396</span>  {</div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span>  <span class="keywordflow">if</span>(lid < 128)</div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>  {</div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>  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>  }</div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span>  }</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 128)</div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>  {</div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>  <span class="keywordflow">if</span>(lid < 64)</div><div class="line"><a name="l00406"></a><span class="lineno"> 406</span>  {</div><div class="line"><a name="l00407"></a><span class="lineno"> 407</span>  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>  }</div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span>  }</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 64)</div><div class="line"><a name="l00412"></a><span class="lineno"> 412</span>  {</div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>  <span class="keywordflow">if</span>(lid < 32)</div><div class="line"><a name="l00414"></a><span class="lineno"> 414</span>  {</div><div class="line"><a name="l00415"></a><span class="lineno"> 415</span>  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>  }</div><div class="line"><a name="l00417"></a><span class="lineno"> 417</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00418"></a><span class="lineno"> 418</span>  }</div><div class="line"><a name="l00419"></a><span class="lineno"> 419</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 32)</div><div class="line"><a name="l00420"></a><span class="lineno"> 420</span>  {</div><div class="line"><a name="l00421"></a><span class="lineno"> 421</span>  <span class="keywordflow">if</span>(lid < 16)</div><div class="line"><a name="l00422"></a><span class="lineno"> 422</span>  {</div><div class="line"><a name="l00423"></a><span class="lineno"> 423</span>  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>  }</div><div class="line"><a name="l00425"></a><span class="lineno"> 425</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00426"></a><span class="lineno"> 426</span>  }</div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 16)</div><div class="line"><a name="l00428"></a><span class="lineno"> 428</span>  {</div><div class="line"><a name="l00429"></a><span class="lineno"> 429</span>  <span class="keywordflow">if</span>(lid < 8)</div><div class="line"><a name="l00430"></a><span class="lineno"> 430</span>  {</div><div class="line"><a name="l00431"></a><span class="lineno"> 431</span>  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>  }</div><div class="line"><a name="l00433"></a><span class="lineno"> 433</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00434"></a><span class="lineno"> 434</span>  }</div><div class="line"><a name="l00435"></a><span class="lineno"> 435</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 8)</div><div class="line"><a name="l00436"></a><span class="lineno"> 436</span>  {</div><div class="line"><a name="l00437"></a><span class="lineno"> 437</span>  <span class="keywordflow">if</span>(lid < 4)</div><div class="line"><a name="l00438"></a><span class="lineno"> 438</span>  {</div><div class="line"><a name="l00439"></a><span class="lineno"> 439</span>  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>  }</div><div class="line"><a name="l00441"></a><span class="lineno"> 441</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00442"></a><span class="lineno"> 442</span>  }</div><div class="line"><a name="l00443"></a><span class="lineno"> 443</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 4)</div><div class="line"><a name="l00444"></a><span class="lineno"> 444</span>  {</div><div class="line"><a name="l00445"></a><span class="lineno"> 445</span>  <span class="keywordflow">if</span>(lid < 2)</div><div class="line"><a name="l00446"></a><span class="lineno"> 446</span>  {</div><div class="line"><a name="l00447"></a><span class="lineno"> 447</span>  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>  }</div><div class="line"><a name="l00449"></a><span class="lineno"> 449</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00450"></a><span class="lineno"> 450</span>  }</div><div class="line"><a name="l00451"></a><span class="lineno"> 451</span>  <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00452"></a><span class="lineno"> 452</span>  {</div><div class="line"><a name="l00453"></a><span class="lineno"> 453</span>  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>  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>  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>  max_local = max_val_vec.s0;</div><div class="line"><a name="l00457"></a><span class="lineno"> 457</span>  }</div><div class="line"><a name="l00458"></a><span class="lineno"> 458</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00459"></a><span class="lineno"> 459</span> </div><div class="line"><a name="l00460"></a><span class="lineno"> 460</span>  <span class="comment">/* Second section */</span></div><div class="line"><a name="l00461"></a><span class="lineno"> 461</span> </div><div class="line"><a name="l00462"></a><span class="lineno"> 462</span>  <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00463"></a><span class="lineno"> 463</span>  <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>  sum1D = 0;</div><div class="line"><a name="l00465"></a><span class="lineno"> 465</span>  <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> </div><div class="line"><a name="l00467"></a><span class="lineno"> 467</span>  <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00468"></a><span class="lineno"> 468</span>  <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a0886942393a3ba0dfefaa7516b159784">for</a>(i = 0; i < width_; i++)</div><div class="line"><a name="l00469"></a><span class="lineno"> 469</span>  {</div><div class="line"><a name="l00470"></a><span class="lineno"> 470</span>  <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>  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>(&<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>  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> <span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00474"></a><span class="lineno"> 474</span>  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> <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> <span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00477"></a><span class="lineno"> 477</span>  <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>  (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>(&<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>  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> <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>  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>  <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>  (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>(&<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> <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>  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>  }</div><div class="line"><a name="l00487"></a><span class="lineno"> 487</span> <span class="preprocessor">#ifdef NON_MULTIPLE_OF_GRID_SIZE</span></div><div class="line"><a name="l00488"></a><span class="lineno"> 488</span>  <span class="comment">//TODO: Optimize the calculation (avoid %).</span></div><div class="line"><a name="l00489"></a><span class="lineno"> 489</span>  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>  <span class="keywordflow">if</span>(lid < boundary_workitems)</div><div class="line"><a name="l00491"></a><span class="lineno"> 491</span>  {</div><div class="line"><a name="l00492"></a><span class="lineno"> 492</span>  <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>  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>(&<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>  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> <span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00496"></a><span class="lineno"> 496</span>  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> <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> <span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00499"></a><span class="lineno"> 499</span>  <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>  (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>(&<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>  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> <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>  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>  <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>  (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>(&<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> <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>  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>  }</div><div class="line"><a name="l00509"></a><span class="lineno"> 509</span> <span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00510"></a><span class="lineno"> 510</span>  <span class="keywordflow">if</span>(boundary_workitems == 0)</div><div class="line"><a name="l00511"></a><span class="lineno"> 511</span>  {</div><div class="line"><a name="l00512"></a><span class="lineno"> 512</span>  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>  i--;</div><div class="line"><a name="l00514"></a><span class="lineno"> 514</span>  }</div><div class="line"><a name="l00515"></a><span class="lineno"> 515</span>  <span class="keywordflow">if</span>(lid == (boundary_workitems - 1))</div><div class="line"><a name="l00516"></a><span class="lineno"> 516</span>  {</div><div class="line"><a name="l00517"></a><span class="lineno"> 517</span>  <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>  <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>  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>(&<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>  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> <span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00522"></a><span class="lineno"> 522</span>  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> <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> <span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00525"></a><span class="lineno"> 525</span>  <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>  (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>(&<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>  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>  <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>  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>) < 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>  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> <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>  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>  <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>  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>) < 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>  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>  <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>  (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>(&<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> <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>  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>  }</div><div class="line"><a name="l00541"></a><span class="lineno"> 541</span> <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> <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>  tmp_local[lid] = sum1D;</div><div class="line"><a name="l00544"></a><span class="lineno"> 544</span> </div><div class="line"><a name="l00545"></a><span class="lineno"> 545</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00546"></a><span class="lineno"> 546</span> </div><div class="line"><a name="l00547"></a><span class="lineno"> 547</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 256)</div><div class="line"><a name="l00548"></a><span class="lineno"> 548</span>  {</div><div class="line"><a name="l00549"></a><span class="lineno"> 549</span>  <span class="keywordflow">if</span>(lid < 128)</div><div class="line"><a name="l00550"></a><span class="lineno"> 550</span>  {</div><div class="line"><a name="l00551"></a><span class="lineno"> 551</span>  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>  }</div><div class="line"><a name="l00553"></a><span class="lineno"> 553</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00554"></a><span class="lineno"> 554</span>  }</div><div class="line"><a name="l00555"></a><span class="lineno"> 555</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 128)</div><div class="line"><a name="l00556"></a><span class="lineno"> 556</span>  {</div><div class="line"><a name="l00557"></a><span class="lineno"> 557</span>  <span class="keywordflow">if</span>(lid < 64)</div><div class="line"><a name="l00558"></a><span class="lineno"> 558</span>  {</div><div class="line"><a name="l00559"></a><span class="lineno"> 559</span>  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>  }</div><div class="line"><a name="l00561"></a><span class="lineno"> 561</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00562"></a><span class="lineno"> 562</span>  }</div><div class="line"><a name="l00563"></a><span class="lineno"> 563</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 64)</div><div class="line"><a name="l00564"></a><span class="lineno"> 564</span>  {</div><div class="line"><a name="l00565"></a><span class="lineno"> 565</span>  <span class="keywordflow">if</span>(lid < 32)</div><div class="line"><a name="l00566"></a><span class="lineno"> 566</span>  {</div><div class="line"><a name="l00567"></a><span class="lineno"> 567</span>  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>  }</div><div class="line"><a name="l00569"></a><span class="lineno"> 569</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00570"></a><span class="lineno"> 570</span>  }</div><div class="line"><a name="l00571"></a><span class="lineno"> 571</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 32)</div><div class="line"><a name="l00572"></a><span class="lineno"> 572</span>  {</div><div class="line"><a name="l00573"></a><span class="lineno"> 573</span>  <span class="keywordflow">if</span>(lid < 16)</div><div class="line"><a name="l00574"></a><span class="lineno"> 574</span>  {</div><div class="line"><a name="l00575"></a><span class="lineno"> 575</span>  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>  }</div><div class="line"><a name="l00577"></a><span class="lineno"> 577</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00578"></a><span class="lineno"> 578</span>  }</div><div class="line"><a name="l00579"></a><span class="lineno"> 579</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 16)</div><div class="line"><a name="l00580"></a><span class="lineno"> 580</span>  {</div><div class="line"><a name="l00581"></a><span class="lineno"> 581</span>  <span class="keywordflow">if</span>(lid < 8)</div><div class="line"><a name="l00582"></a><span class="lineno"> 582</span>  {</div><div class="line"><a name="l00583"></a><span class="lineno"> 583</span>  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>  }</div><div class="line"><a name="l00585"></a><span class="lineno"> 585</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00586"></a><span class="lineno"> 586</span>  }</div><div class="line"><a name="l00587"></a><span class="lineno"> 587</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 8)</div><div class="line"><a name="l00588"></a><span class="lineno"> 588</span>  {</div><div class="line"><a name="l00589"></a><span class="lineno"> 589</span>  <span class="keywordflow">if</span>(lid < 4)</div><div class="line"><a name="l00590"></a><span class="lineno"> 590</span>  {</div><div class="line"><a name="l00591"></a><span class="lineno"> 591</span>  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>  }</div><div class="line"><a name="l00593"></a><span class="lineno"> 593</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00594"></a><span class="lineno"> 594</span>  }</div><div class="line"><a name="l00595"></a><span class="lineno"> 595</span>  <span class="keywordflow">if</span>(<a class="code" href="softmax__layer_8cl.xhtml#a08246606c233e7785a497c09672f366f">GRID_SIZE</a> >= 4)</div><div class="line"><a name="l00596"></a><span class="lineno"> 596</span>  {</div><div class="line"><a name="l00597"></a><span class="lineno"> 597</span>  <span class="keywordflow">if</span>(lid < 2)</div><div class="line"><a name="l00598"></a><span class="lineno"> 598</span>  {</div><div class="line"><a name="l00599"></a><span class="lineno"> 599</span>  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>  }</div><div class="line"><a name="l00601"></a><span class="lineno"> 601</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00602"></a><span class="lineno"> 602</span>  }</div><div class="line"><a name="l00603"></a><span class="lineno"> 603</span>  <span class="keywordflow">if</span>(lid == 0)</div><div class="line"><a name="l00604"></a><span class="lineno"> 604</span>  {</div><div class="line"><a name="l00605"></a><span class="lineno"> 605</span>  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>  <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00607"></a><span class="lineno"> 607</span>  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>  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>  *((__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>  }</div><div class="line"><a name="l00611"></a><span class="lineno"> 611</span> }</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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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#l00066">reduction_operation.cl:66</a></div></div> |
| <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> |
| <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> |
| <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< _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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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 & src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div> |
| <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> |
| </div><!-- fragment --> |
| <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#l00066">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> |
| |
| </div> |
| </div> |
| <a id="a9d74601bcabbb2f14bcf52385ad666dc"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a9d74601bcabbb2f14bcf52385ad666dc">◆ </a></span>softmax_layer_max_shift_exp_sum_serial()</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__kernel void softmax_layer_max_shift_exp_sum_serial </td> |
| <td>(</td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>src_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>maxo_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>maxo_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>dst_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>sum_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>width</em> </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td></td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| <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> |
| <dd> |
| In case the input is not a multiple of VECTOR_SIZE (2,4,8,16) -DNON_MULTIPLE_OF_VECTOR_SIZE must be passed. </dd> |
| <dd> |
| Beta can be optionally passed at compile time using -DBETA (by default, it is 1.0).</dd></dl> |
| <dl class="params"><dt>Parameters</dt><dd> |
| <table class="params"> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <tr><td class="paramdir">[in]</td><td class="paramname">width</td><td>Input image width </td></tr> |
| </table> |
| </dd> |
| </dl> |
| |
| <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> |
| <div class="fragment"><div class="line"><a name="l00168"></a><span class="lineno"> 168</span> {</div><div class="line"><a name="l00169"></a><span class="lineno"> 169</span>  <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>  <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>  <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>  <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> </div><div class="line"><a name="l00174"></a><span class="lineno"> 174</span> <span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00175"></a><span class="lineno"> 175</span>  <span class="comment">// Initialize beta</span></div><div class="line"><a name="l00176"></a><span class="lineno"> 176</span>  <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>  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> <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> </div><div class="line"><a name="l00180"></a><span class="lineno"> 180</span>  <span class="comment">// Initialize local maximum</span></div><div class="line"><a name="l00181"></a><span class="lineno"> 181</span>  <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>  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> </div><div class="line"><a name="l00184"></a><span class="lineno"> 184</span>  <span class="comment">// Calculate max of row</span></div><div class="line"><a name="l00185"></a><span class="lineno"> 185</span>  <span class="keyword">const</span> uint width_ = width >> <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>  <span class="keywordflow">for</span>(uint i = 0; i < width_; i++)</div><div class="line"><a name="l00187"></a><span class="lineno"> 187</span>  {</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>  <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>  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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i << <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>  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>  }</div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span> </div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span> <span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>  <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>  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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, width_ << <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>  <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>  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_ << <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) < 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>  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> <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> </div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span>  <span class="comment">// Perform max reduction</span></div><div class="line"><a name="l00202"></a><span class="lineno"> 202</span> <span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>  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> <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> <span class="preprocessor">#if VECTOR_SIZE >= 8</span></div><div class="line"><a name="l00206"></a><span class="lineno"> 206</span>  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> <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> <span class="preprocessor">#if VECTOR_SIZE >= 4</span></div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span>  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> <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>  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>  <span class="comment">// Store result</span></div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>  *((__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> </div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>  <span class="comment">/* Second section */</span></div><div class="line"><a name="l00216"></a><span class="lineno"> 216</span> </div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>  <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>  <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>(&maxo, 0, 0));</div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span> </div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span>  <span class="comment">// Set sum vector</span></div><div class="line"><a name="l00221"></a><span class="lineno"> 221</span>  <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>  sum1D = 0;</div><div class="line"><a name="l00223"></a><span class="lineno"> 223</span> </div><div class="line"><a name="l00224"></a><span class="lineno"> 224</span>  <span class="comment">// Shift values, exp and sum</span></div><div class="line"><a name="l00225"></a><span class="lineno"> 225</span>  <span class="keywordflow">for</span>(uint i = 0; i < width_; i++)</div><div class="line"><a name="l00226"></a><span class="lineno"> 226</span>  {</div><div class="line"><a name="l00227"></a><span class="lineno"> 227</span>  <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>  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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, i << <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>  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> <span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00231"></a><span class="lineno"> 231</span>  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> <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> <span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00234"></a><span class="lineno"> 234</span>  <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>  (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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i << <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>  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> <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>  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>  <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>  (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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, i << <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> <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>  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>  }</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span> </div><div class="line"><a name="l00245"></a><span class="lineno"> 245</span> <span class="preprocessor">#ifdef NON_MULTIPLE_OF_VECTOR_SIZE</span></div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span>  <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>  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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a989ab3e96426615bb98e04e0235088ca">src</a>, width_ << <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>  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> <span class="preprocessor">#ifdef BETA</span></div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>  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> <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> <span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span>  <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>  (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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, width_ << <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>  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>  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_ << <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) < 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>  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> <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>  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>  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_ << <a class="code" href="softmax__layer_8cl.xhtml#a372393c380805985b813dbb16d589a64">LOG_VECTOR_SIZE</a>) + <a class="code" href="softmax__layer_8cl.xhtml#aa1dd94b8d98f1c6d790bdf0fc5de29e9">idx__</a>) < 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>  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>  <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>  (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>(&<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>, width_ << <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> <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>  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> <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> </div><div class="line"><a name="l00268"></a><span class="lineno"> 268</span>  <span class="comment">// Perform sum reduction</span></div><div class="line"><a name="l00269"></a><span class="lineno"> 269</span> <span class="preprocessor">#if VECTOR_SIZE == 16</span></div><div class="line"><a name="l00270"></a><span class="lineno"> 270</span>  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> <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> <span class="preprocessor">#if VECTOR_SIZE >= 8</span></div><div class="line"><a name="l00273"></a><span class="lineno"> 273</span>  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> <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> <span class="preprocessor">#if VECTOR_SIZE >= 4</span></div><div class="line"><a name="l00276"></a><span class="lineno"> 276</span>  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> <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>  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> </div><div class="line"><a name="l00280"></a><span class="lineno"> 280</span>  <span class="comment">// Calculate and store result</span></div><div class="line"><a name="l00281"></a><span class="lineno"> 281</span>  *((__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> }</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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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#l00066">reduction_operation.cl:66</a></div></div> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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 & src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div> |
| <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> |
| </div><!-- fragment --> |
| <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#l00066">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> |
| |
| </div> |
| </div> |
| <a id="ac4247ac0991e85965b7ded764e78f12c"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#ac4247ac0991e85965b7ded764e78f12c">◆ </a></span>softmax_layer_norm()</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__kernel void softmax_layer_norm </td> |
| <td>(</td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>src_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>sum_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">__global uchar * </td> |
| <td class="paramname"><em>dst_ptr</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_x</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_y</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_stride_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_step_z</em>, </td> |
| </tr> |
| <tr> |
| <td class="paramkey"></td> |
| <td></td> |
| <td class="paramtype">uint </td> |
| <td class="paramname"><em>dst_offset_first_element_in_bytes</em> </td> |
| </tr> |
| <tr> |
| <td></td> |
| <td>)</td> |
| <td></td><td></td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <p>Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel. </p> |
| <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> |
| <dl class="params"><dt>Parameters</dt><dd> |
| <table class="params"> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| </table> |
| </dd> |
| </dl> |
| |
| <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> |
| <div class="fragment"><div class="line"><a name="l00105"></a><span class="lineno"> 105</span> {</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span>  <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>  <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>  <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> </div><div class="line"><a name="l00110"></a><span class="lineno"> 110</span>  <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>  <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>(&<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>  <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>  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>(&<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> <span class="preprocessor">#ifdef LOG_SOFTMAX</span></div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span>  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>(&<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> <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>  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>(&<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> <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> }</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> |
| <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> |
| <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> |
| <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#l00066">reduction_operation.cl:66</a></div></div> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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> |
| <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 & src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div> |
| <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> |
| </div><!-- fragment --> |
| <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#l00066">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> |
| |
| </div> |
| </div> |
| <h2 class="groupheader">Variable Documentation</h2> |
| <a id="a0712735973f172ac9efc7d48a31e47ad"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a0712735973f172ac9efc7d48a31e47ad">◆ </a></span>idx16</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="a4884a666a1e93fbf8c27bd7d2da3c8bb"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a4884a666a1e93fbf8c27bd7d2da3c8bb">◆ </a></span>idx4</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__constant uint4 idx4 = (uint4)(0, 1, 2, 3)</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| <p class="reference">Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00325">softmax_layer_max_shift_exp_sum_parallel()</a>.</p> |
| |
| </div> |
| </div> |
| <a id="aa1dd94b8d98f1c6d790bdf0fc5de29e9"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#aa1dd94b8d98f1c6d790bdf0fc5de29e9">◆ </a></span>idx__</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__constant uint16 idx__ = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| <p class="reference">Referenced by <a class="el" href="softmax__layer_8cl_source.xhtml#l00162">softmax_layer_max_shift_exp_sum_serial()</a>.</p> |
| |
| </div> |
| </div> |
| <a id="a538b4b63f40e7b12891774e03a4f0dec"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#a538b4b63f40e7b12891774e03a4f0dec">◆ </a></span>type_min</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( -FLT_MAX )</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| </div> |
| </div> |
| <a id="af7a08044d0e491a0ee1520a24a107a2b"></a> |
| <h2 class="memtitle"><span class="permalink"><a href="#af7a08044d0e491a0ee1520a24a107a2b">◆ </a></span>type_min_</h2> |
| |
| <div class="memitem"> |
| <div class="memproto"> |
| <table class="memname"> |
| <tr> |
| <td class="memname">__constant DATA_TYPE16 type_min_ = ( DATA_TYPE16 )( -FLT_MAX )</td> |
| </tr> |
| </table> |
| </div><div class="memdoc"> |
| |
| <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> |
| |
| <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> |
| |
| </div> |
| </div> |
| </div><!-- contents --> |
| </div><!-- doc-content --> |
| <!-- start footer part --> |
| <div id="nav-path" class="navpath"><!-- id is needed for treeview function! --> |
| <ul> |
| <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> |
| <li class="footer">Generated on Thu Mar 5 2020 16:07:07 for Compute Library by |
| <a href="http://www.doxygen.org/index.html"> |
| <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.15 </li> |
| </ul> |
| </div> |
| </body> |
| </html> |
