Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 52ba29e936b8e711e8acdfe819e36f884d4f3fe1 / . / documentation / normalization__layer_8cl.xhtml
blob: 56e135ae9befd1e0844bac0acee3db101dab17b8 [file] [log] [blame]
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1<!-- HTML header for doxygen 1.8.9.1-->
2<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
3<html xmlns="http://www.w3.org/1999/xhtml">
4<head>
5<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
6<meta http-equiv="X-UA-Compatible" content="IE=9"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]7<meta name="generator" content="Doxygen 1.8.11"/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]8<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]9<title>Compute Library: src/core/CL/cl_kernels/normalization_layer.cl File Reference</title>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]10<link href="tabs.css" rel="stylesheet" type="text/css"/>
11<script type="text/javascript" src="jquery.js"></script>
12<script type="text/javascript" src="dynsections.js"></script>
13<link href="navtree.css" rel="stylesheet" type="text/css"/>
14<script type="text/javascript" src="resize.js"></script>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]15<script type="text/javascript" src="navtreedata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]16<script type="text/javascript" src="navtree.js"></script>
17<script type="text/javascript">
18 $(document).ready(initResizable);
19 $(window).load(resizeHeight);
20</script>
21<link href="search/search.css" rel="stylesheet" type="text/css"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]22<script type="text/javascript" src="search/searchdata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]23<script type="text/javascript" src="search/search.js"></script>
24<script type="text/javascript">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]25 $(document).ready(function() { init_search(); });
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]26</script>
27<script type="text/x-mathjax-config">
28 MathJax.Hub.Config({
29 extensions: ["tex2jax.js"],
30 jax: ["input/TeX","output/HTML-CSS"],
31});
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]32</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]33<link href="doxygen.css" rel="stylesheet" type="text/css" />
34</head>
35<body>
36<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
37<div id="titlearea">
38<table cellspacing="0" cellpadding="0">
39 <tbody>
40 <tr style="height: 56px;">
41 <td style="padding-left: 0.5em;">
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]42 <div id="projectname">Compute Library
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]43 &#160;<span id="projectnumber">18.08</span>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]44 </div>
45 </td>
46 </tr>
47 </tbody>
48</table>
49</div>
50<!-- end header part -->
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]51<!-- Generated by Doxygen 1.8.11 -->
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]52<script type="text/javascript">
53var searchBox = new SearchBox("searchBox", "search",false,'Search');
54</script>
55 <div id="navrow1" class="tabs">
56 <ul class="tablist">
57 <li><a href="index.xhtml"><span>Main&#160;Page</span></a></li>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]58 <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]59 <li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
60 <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
61 <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
62 <li>
63 <div id="MSearchBox" class="MSearchBoxInactive">
64 <span class="left">
65 <img id="MSearchSelect" src="search/mag_sel.png"
66 onmouseover="return searchBox.OnSearchSelectShow()"
67 onmouseout="return searchBox.OnSearchSelectHide()"
68 alt=""/>
69 <input type="text" id="MSearchField" value="Search" accesskey="S"
70 onfocus="searchBox.OnSearchFieldFocus(true)"
71 onblur="searchBox.OnSearchFieldFocus(false)"
72 onkeyup="searchBox.OnSearchFieldChange(event)"/>
73 </span><span class="right">
74 <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
75 </span>
76 </div>
77 </li>
78 </ul>
79 </div>
80 <div id="navrow2" class="tabs2">
81 <ul class="tablist">
82 <li><a href="files.xhtml"><span>File&#160;List</span></a></li>
83 <li><a href="globals.xhtml"><span>Globals</span></a></li>
84 </ul>
85 </div>
86</div><!-- top -->
87<div id="side-nav" class="ui-resizable side-nav-resizable">
88 <div id="nav-tree">
89 <div id="nav-tree-contents">
90 <div id="nav-sync" class="sync"></div>
91 </div>
92 </div>
93 <div id="splitbar" style="-moz-user-select:none;"
94 class="ui-resizable-handle">
95 </div>
96</div>
97<script type="text/javascript">
98$(document).ready(function(){initNavTree('normalization__layer_8cl.xhtml','');});
99</script>
100<div id="doc-content">
101<!-- window showing the filter options -->
102<div id="MSearchSelectWindow"
103 onmouseover="return searchBox.OnSearchSelectShow()"
104 onmouseout="return searchBox.OnSearchSelectHide()"
105 onkeydown="return searchBox.OnSearchSelectKey(event)">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]106</div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]107
108<!-- iframe showing the search results (closed by default) -->
109<div id="MSearchResultsWindow">
110<iframe src="javascript:void(0)" frameborder="0"
111 name="MSearchResults" id="MSearchResults">
112</iframe>
113</div>
114
115<div class="header">
116 <div class="summary">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]117<a href="#define-members">Macros</a> &#124;
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]118<a href="#func-members">Functions</a> </div>
119 <div class="headertitle">
120<div class="title">normalization_layer.cl File Reference</div> </div>
121</div><!--header-->
122<div class="contents">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]123<div class="textblock"><code>#include &quot;<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]124</div>
125<p><a href="normalization__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
126<table class="memberdecls">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]127<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
128Macros</h2></td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]129<tr class="memitem:ac0bf948069436897cf005ea45ac73a4f"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#ac0bf948069436897cf005ea45ac73a4f">MUL_OP</a>(x, y)&#160;&#160;&#160;((x) * (y))</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]130<tr class="separator:ac0bf948069436897cf005ea45ac73a4f"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]131<tr class="memitem:a334426d638df3b4108c366e240cf5504"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#a334426d638df3b4108c366e240cf5504">ADD_OP</a>(x, y)&#160;&#160;&#160;((x) + (y))</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]132<tr class="separator:a334426d638df3b4108c366e240cf5504"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]133<tr class="memitem:a5db17889d824975fefb2ce2f4690637f"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a>(x, y)&#160;&#160;&#160;((x) / (y))</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]134<tr class="separator:a5db17889d824975fefb2ce2f4690637f"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]135<tr class="memitem:a09724da1a9e5ba54260905829ec27db4"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#a09724da1a9e5ba54260905829ec27db4">POW_OP</a>(x, y)&#160;&#160;&#160;pow((x), (y))</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]136<tr class="separator:a09724da1a9e5ba54260905829ec27db4"><td class="memSeparator" colspan="2">&#160;</td></tr>
137<tr class="memitem:a107d847044e677b01e9bd3d5251b39d9"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a>(a)&#160;&#160;&#160;(a)</td></tr>
138<tr class="separator:a107d847044e677b01e9bd3d5251b39d9"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]139<tr class="memitem:a10a9bf1a9fce27664d5f91844ff3f691"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#a10a9bf1a9fce27664d5f91844ff3f691">LOAD_OP</a>(<a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)&#160;&#160;&#160;vload4(<a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]140<tr class="separator:a10a9bf1a9fce27664d5f91844ff3f691"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]141<tr class="memitem:ac59caf44b60378c02032aab994ad5ac1"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#ac59caf44b60378c02032aab994ad5ac1">STORE_OP</a>(data, <a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)&#160;&#160;&#160;vstore4(data, <a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]142<tr class="separator:ac59caf44b60378c02032aab994ad5ac1"><td class="memSeparator" colspan="2">&#160;</td></tr>
143</table><table class="memberdecls">
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]144<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
145Functions</h2></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]146<tr class="memitem:a7db22b7ddf8f433e02ede3bd9c56d9d9"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#a7db22b7ddf8f433e02ede3bd9c56d9d9">normalization_layer_cross_map</a> (__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]147<tr class="memdesc:a7db22b7ddf8f433e02ede3bd9c56d9d9"><td class="mdescLeft">&#160;</td><td class="mdescRight">Apply cross-map normalization. <a href="#a7db22b7ddf8f433e02ede3bd9c56d9d9">More...</a><br /></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]148<tr class="separator:a7db22b7ddf8f433e02ede3bd9c56d9d9"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]149<tr class="memitem:a98051a32b8ae24598346306c1b131b2c"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="normalization__layer_8cl.xhtml#a98051a32b8ae24598346306c1b131b2c">normalization_layer_in_map</a> (__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_offset_first_element_in_bytes)</td></tr>
150<tr class="memdesc:a98051a32b8ae24598346306c1b131b2c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Apply in-map normalization. <a href="#a98051a32b8ae24598346306c1b131b2c">More...</a><br /></td></tr>
151<tr class="separator:a98051a32b8ae24598346306c1b131b2c"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]152</table>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]153<h2 class="groupheader">Macro Definition Documentation</h2>
154<a class="anchor" id="a334426d638df3b4108c366e240cf5504"></a>
155<div class="memitem">
156<div class="memproto">
157 <table class="memname">
158 <tr>
159 <td class="memname">#define ADD_OP</td>
160 <td>(</td>
161 <td class="paramtype">&#160;</td>
162 <td class="paramname">x, </td>
163 </tr>
164 <tr>
165 <td class="paramkey"></td>
166 <td></td>
167 <td class="paramtype">&#160;</td>
168 <td class="paramname">y&#160;</td>
169 </tr>
170 <tr>
171 <td></td>
172 <td>)</td>
173 <td></td><td>&#160;&#160;&#160;((x) + (y))</td>
174 </tr>
175 </table>
176</div><div class="memdoc">
177
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]178<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00027">27</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]179
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]180<p>Referenced by <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">normalization_layer_cross_map()</a>, and <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">normalization_layer_in_map()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]181
182</div>
183</div>
184<a class="anchor" id="a5db17889d824975fefb2ce2f4690637f"></a>
185<div class="memitem">
186<div class="memproto">
187 <table class="memname">
188 <tr>
189 <td class="memname">#define DIV_OP</td>
190 <td>(</td>
191 <td class="paramtype">&#160;</td>
192 <td class="paramname">x, </td>
193 </tr>
194 <tr>
195 <td class="paramkey"></td>
196 <td></td>
197 <td class="paramtype">&#160;</td>
198 <td class="paramname">y&#160;</td>
199 </tr>
200 <tr>
201 <td></td>
202 <td>)</td>
203 <td></td><td>&#160;&#160;&#160;((x) / (y))</td>
204 </tr>
205 </table>
206</div><div class="memdoc">
207
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]208<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00028">28</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]209
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]210<p>Referenced by <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">normalization_layer_cross_map()</a>, and <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">normalization_layer_in_map()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]211
212</div>
213</div>
214<a class="anchor" id="a10a9bf1a9fce27664d5f91844ff3f691"></a>
215<div class="memitem">
216<div class="memproto">
217 <table class="memname">
218 <tr>
219 <td class="memname">#define LOAD_OP</td>
220 <td>(</td>
221 <td class="paramtype">&#160;</td>
222 <td class="paramname"><a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, </td>
223 </tr>
224 <tr>
225 <td class="paramkey"></td>
226 <td></td>
227 <td class="paramtype">&#160;</td>
228 <td class="paramname">ptr&#160;</td>
229 </tr>
230 <tr>
231 <td></td>
232 <td>)</td>
233 <td></td><td>&#160;&#160;&#160;vload4(<a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)</td>
234 </tr>
235 </table>
236</div><div class="memdoc">
237
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]238<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00032">32</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]239
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]240<p>Referenced by <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">normalization_layer_cross_map()</a>, and <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">normalization_layer_in_map()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]241
242</div>
243</div>
244<a class="anchor" id="ac0bf948069436897cf005ea45ac73a4f"></a>
245<div class="memitem">
246<div class="memproto">
247 <table class="memname">
248 <tr>
249 <td class="memname">#define MUL_OP</td>
250 <td>(</td>
251 <td class="paramtype">&#160;</td>
252 <td class="paramname">x, </td>
253 </tr>
254 <tr>
255 <td class="paramkey"></td>
256 <td></td>
257 <td class="paramtype">&#160;</td>
258 <td class="paramname">y&#160;</td>
259 </tr>
260 <tr>
261 <td></td>
262 <td>)</td>
263 <td></td><td>&#160;&#160;&#160;((x) * (y))</td>
264 </tr>
265 </table>
266</div><div class="memdoc">
267
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]268<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00026">26</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]269
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]270<p>Referenced by <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">normalization_layer_cross_map()</a>, and <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">normalization_layer_in_map()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]271
272</div>
273</div>
274<a class="anchor" id="a09724da1a9e5ba54260905829ec27db4"></a>
275<div class="memitem">
276<div class="memproto">
277 <table class="memname">
278 <tr>
279 <td class="memname">#define POW_OP</td>
280 <td>(</td>
281 <td class="paramtype">&#160;</td>
282 <td class="paramname">x, </td>
283 </tr>
284 <tr>
285 <td class="paramkey"></td>
286 <td></td>
287 <td class="paramtype">&#160;</td>
288 <td class="paramname">y&#160;</td>
289 </tr>
290 <tr>
291 <td></td>
292 <td>)</td>
293 <td></td><td>&#160;&#160;&#160;pow((x), (y))</td>
294 </tr>
295 </table>
296</div><div class="memdoc">
297
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]298<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00029">29</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]299
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]300<p>Referenced by <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">normalization_layer_cross_map()</a>, and <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">normalization_layer_in_map()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]301
302</div>
303</div>
304<a class="anchor" id="a107d847044e677b01e9bd3d5251b39d9"></a>
305<div class="memitem">
306<div class="memproto">
307 <table class="memname">
308 <tr>
309 <td class="memname">#define SQCVT_SAT</td>
310 <td>(</td>
311 <td class="paramtype">&#160;</td>
312 <td class="paramname">a</td><td>)</td>
313 <td>&#160;&#160;&#160;(a)</td>
314 </tr>
315 </table>
316</div><div class="memdoc">
317
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]318<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00030">30</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]319
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]320<p>Referenced by <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">normalization_layer_cross_map()</a>, and <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">normalization_layer_in_map()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]321
322</div>
323</div>
324<a class="anchor" id="ac59caf44b60378c02032aab994ad5ac1"></a>
325<div class="memitem">
326<div class="memproto">
327 <table class="memname">
328 <tr>
329 <td class="memname">#define STORE_OP</td>
330 <td>(</td>
331 <td class="paramtype">&#160;</td>
332 <td class="paramname">data, </td>
333 </tr>
334 <tr>
335 <td class="paramkey"></td>
336 <td></td>
337 <td class="paramtype">&#160;</td>
338 <td class="paramname"><a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, </td>
339 </tr>
340 <tr>
341 <td class="paramkey"></td>
342 <td></td>
343 <td class="paramtype">&#160;</td>
344 <td class="paramname">ptr&#160;</td>
345 </tr>
346 <tr>
347 <td></td>
348 <td>)</td>
349 <td></td><td>&#160;&#160;&#160;vstore4(data, <a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>, ptr)</td>
350 </tr>
351 </table>
352</div><div class="memdoc">
353
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]354<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00033">33</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]355
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]356<p>Referenced by <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">normalization_layer_cross_map()</a>, and <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">normalization_layer_in_map()</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]357
358</div>
359</div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]360<h2 class="groupheader">Function Documentation</h2>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]361<a class="anchor" id="a7db22b7ddf8f433e02ede3bd9c56d9d9"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]362<div class="memitem">
363<div class="memproto">
364 <table class="memname">
365 <tr>
366 <td class="memname">__kernel void normalization_layer_cross_map </td>
367 <td>(</td>
368 <td class="paramtype">__global uchar *&#160;</td>
369 <td class="paramname"><em>input_ptr</em>, </td>
370 </tr>
371 <tr>
372 <td class="paramkey"></td>
373 <td></td>
374 <td class="paramtype">uint&#160;</td>
375 <td class="paramname"><em>input_stride_x</em>, </td>
376 </tr>
377 <tr>
378 <td class="paramkey"></td>
379 <td></td>
380 <td class="paramtype">uint&#160;</td>
381 <td class="paramname"><em>input_step_x</em>, </td>
382 </tr>
383 <tr>
384 <td class="paramkey"></td>
385 <td></td>
386 <td class="paramtype">uint&#160;</td>
387 <td class="paramname"><em>input_stride_y</em>, </td>
388 </tr>
389 <tr>
390 <td class="paramkey"></td>
391 <td></td>
392 <td class="paramtype">uint&#160;</td>
393 <td class="paramname"><em>input_step_y</em>, </td>
394 </tr>
395 <tr>
396 <td class="paramkey"></td>
397 <td></td>
398 <td class="paramtype">uint&#160;</td>
399 <td class="paramname"><em>input_stride_z</em>, </td>
400 </tr>
401 <tr>
402 <td class="paramkey"></td>
403 <td></td>
404 <td class="paramtype">uint&#160;</td>
405 <td class="paramname"><em>input_step_z</em>, </td>
406 </tr>
407 <tr>
408 <td class="paramkey"></td>
409 <td></td>
410 <td class="paramtype">uint&#160;</td>
411 <td class="paramname"><em>input_offset_first_element_in_bytes</em>, </td>
412 </tr>
413 <tr>
414 <td class="paramkey"></td>
415 <td></td>
416 <td class="paramtype">__global uchar *&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]417 <td class="paramname"><em>output_ptr</em>, </td>
418 </tr>
419 <tr>
420 <td class="paramkey"></td>
421 <td></td>
422 <td class="paramtype">uint&#160;</td>
423 <td class="paramname"><em>output_stride_x</em>, </td>
424 </tr>
425 <tr>
426 <td class="paramkey"></td>
427 <td></td>
428 <td class="paramtype">uint&#160;</td>
429 <td class="paramname"><em>output_step_x</em>, </td>
430 </tr>
431 <tr>
432 <td class="paramkey"></td>
433 <td></td>
434 <td class="paramtype">uint&#160;</td>
435 <td class="paramname"><em>output_stride_y</em>, </td>
436 </tr>
437 <tr>
438 <td class="paramkey"></td>
439 <td></td>
440 <td class="paramtype">uint&#160;</td>
441 <td class="paramname"><em>output_step_y</em>, </td>
442 </tr>
443 <tr>
444 <td class="paramkey"></td>
445 <td></td>
446 <td class="paramtype">uint&#160;</td>
447 <td class="paramname"><em>output_stride_z</em>, </td>
448 </tr>
449 <tr>
450 <td class="paramkey"></td>
451 <td></td>
452 <td class="paramtype">uint&#160;</td>
453 <td class="paramname"><em>output_step_z</em>, </td>
454 </tr>
455 <tr>
456 <td class="paramkey"></td>
457 <td></td>
458 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]459 <td class="paramname"><em>output_offset_first_element_in_bytes</em>&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]460 </tr>
461 <tr>
462 <td></td>
463 <td>)</td>
464 <td></td><td></td>
465 </tr>
466 </table>
467</div><div class="memdoc">
468
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]469<p>Apply cross-map normalization. </p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]470<dl class="section note"><dt>Note</dt><dd>Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short </dd>
471<dd>
472<a class="el" href="struct_vector.xhtml" title="Structure to hold Vector information. ">Vector</a> size should be given as a preprocessor argument using -DVEC_SIZE=size, e.g. -DVEC_SIZE=16 </dd>
473<dd>
474The radius should be given as a preprocessor argument using -DRADIUS=size. e.g. -DRADIUS=5 </dd>
475<dd>
476The number of slices should be given as a preprocessor argument using -DNUM_SLICES=size. e.g. -DNUM_SLICES=192 </dd>
477<dd>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]478Scaling coefficient (= alpha/norm_size), beta and kappa need to be passed at compile time using -DCOEFF, -DALPHA and -DKAPPA</dd></dl>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]479<dl class="params"><dt>Parameters</dt><dd>
480 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]481 <tr><td class="paramdir">[in]</td><td class="paramname">input_ptr</td><td>Pointer to the first source tensor. Supported data types: F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]482 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_x</td><td>Stride of the first source tensor in X dimension (in bytes) </td></tr>
483 <tr><td class="paramdir">[in]</td><td class="paramname">input_step_x</td><td>input_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
484 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_y</td><td>Stride of the first source tensor in Y dimension (in bytes) </td></tr>
485 <tr><td class="paramdir">[in]</td><td class="paramname">input_step_y</td><td>input_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
486 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_z</td><td>Stride of the first source tensor in Z dimension (in bytes) </td></tr>
487 <tr><td class="paramdir">[in]</td><td class="paramname">input_step_z</td><td>input_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
488 <tr><td class="paramdir">[in]</td><td class="paramname">input_offset_first_element_in_bytes</td><td>The offset of the first element in the first source tensor </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]489 <tr><td class="paramdir">[out]</td><td class="paramname">output_ptr</td><td>Pointer to the destination tensor. Supported data types: same as <code>input_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]490 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_x</td><td>Stride of the destination tensor in X dimension (in bytes) </td></tr>
491 <tr><td class="paramdir">[in]</td><td class="paramname">output_step_x</td><td>output_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
492 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_y</td><td>Stride of the destination tensor in Y dimension (in bytes) </td></tr>
493 <tr><td class="paramdir">[in]</td><td class="paramname">output_step_y</td><td>output_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
494 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
495 <tr><td class="paramdir">[in]</td><td class="paramname">output_step_z</td><td>output_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
496 <tr><td class="paramdir">[in]</td><td class="paramname">output_offset_first_element_in_bytes</td><td>The offset of the first element in the destination tensor </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]497 </table>
498 </dd>
499</dl>
500
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]501<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00060">60</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]502
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]503<p>References <a class="el" href="normalization__layer_8cl_source.xhtml#l00027">ADD_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00127">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00028">DIV_OP</a>, <a class="el" href="_c_l_2_winograd_8cpp_source.xhtml#l00241">arm_compute::test::validation::in</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00032">LOAD_OP</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00026">MUL_OP</a>, <a class="el" href="_c_l_2_winograd_8cpp_source.xhtml#l00242">arm_compute::test::validation::out</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00029">POW_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00161">Tensor3D::ptr</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00030">SQCVT_SAT</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00033">STORE_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00323">tensor3D_offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00056">VEC_DATA_TYPE</a>.</p>
504<div class="fragment"><div class="line"><a name="l00062"></a><span class="lineno"> 62</span>&#160;{</div><div class="line"><a name="l00063"></a><span class="lineno"> 63</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a> = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(input);</div><div class="line"><a name="l00064"></a><span class="lineno"> 64</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af8a8a0625e7981212a0af48deb9d2a09">out</a> = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(output);</div><div class="line"><a name="l00065"></a><span class="lineno"> 65</span>&#160;</div><div class="line"><a name="l00066"></a><span class="lineno"> 66</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, VEC_SIZE)</div><div class="line"><a name="l00067"></a><span class="lineno"> 67</span>&#160; acc = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, VEC_SIZE))0;</div><div class="line"><a name="l00068"></a><span class="lineno"> 68</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00069"></a><span class="lineno"> 69</span>&#160; coeff_v = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE))<a class="code" href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a>(COEFF);</div><div class="line"><a name="l00070"></a><span class="lineno"> 70</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00071"></a><span class="lineno"> 71</span>&#160; beta_v = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE))<a class="code" href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a>(BETA);</div><div class="line"><a name="l00072"></a><span class="lineno"> 72</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00073"></a><span class="lineno"> 73</span>&#160; kappa_v = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE))<a class="code" href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a>(KAPPA);</div><div class="line"><a name="l00074"></a><span class="lineno"> 74</span>&#160;</div><div class="line"><a name="l00075"></a><span class="lineno"> 75</span>&#160; const <span class="keywordtype">int</span> current_slice = get_global_id(2);</div><div class="line"><a name="l00076"></a><span class="lineno"> 76</span>&#160; const <span class="keywordtype">int</span> left_slice = max(-(<span class="keywordtype">int</span>)RADIUS, -current_slice);</div><div class="line"><a name="l00077"></a><span class="lineno"> 77</span>&#160; const <span class="keywordtype">int</span> right_slice = min((<span class="keywordtype">int</span>)RADIUS, (<span class="keywordtype">int</span>)NUM_SLICES - 1 - current_slice);</div><div class="line"><a name="l00078"></a><span class="lineno"> 78</span>&#160;</div><div class="line"><a name="l00079"></a><span class="lineno"> 79</span>&#160; for(<span class="keywordtype">int</span> i = left_slice; i &lt;= right_slice; i++)</div><div class="line"><a name="l00080"></a><span class="lineno"> 80</span>&#160; {</div><div class="line"><a name="l00081"></a><span class="lineno"> 81</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00082"></a><span class="lineno"> 82</span>&#160; values = <a class="code" href="normalization__layer_8cl.xhtml#a10a9bf1a9fce27664d5f91844ff3f691">LOAD_OP</a>(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;in, 0, 0, i));</div><div class="line"><a name="l00083"></a><span class="lineno"> 83</span>&#160; acc = <a class="code" href="normalization__layer_8cl.xhtml#a334426d638df3b4108c366e240cf5504">ADD_OP</a>(acc, <a class="code" href="normalization__layer_8cl.xhtml#ac0bf948069436897cf005ea45ac73a4f">MUL_OP</a>(values, values));</div><div class="line"><a name="l00084"></a><span class="lineno"> 84</span>&#160; }</div><div class="line"><a name="l00085"></a><span class="lineno"> 85</span>&#160;</div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span>&#160; acc = <a class="code" href="normalization__layer_8cl.xhtml#a334426d638df3b4108c366e240cf5504">ADD_OP</a>(<a class="code" href="normalization__layer_8cl.xhtml#ac0bf948069436897cf005ea45ac73a4f">MUL_OP</a>(acc, coeff_v), kappa_v);</div><div class="line"><a name="l00087"></a><span class="lineno"> 87</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span>&#160; normalized = <a class="code" href="normalization__layer_8cl.xhtml#a09724da1a9e5ba54260905829ec27db4">POW_OP</a>(acc, beta_v);</div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>&#160; normalized_pixel = <a class="code" href="normalization__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a>(LOAD_OP(0, (__global DATA_TYPE *)in.ptr), normalized);</div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span>&#160;</div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span>&#160; <a class="code" href="normalization__layer_8cl.xhtml#ac59caf44b60378c02032aab994ad5ac1">STORE_OP</a>(normalized_pixel, 0, (__global DATA_TYPE *)out.ptr);</div><div class="line"><a name="l00093"></a><span class="lineno"> 93</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00127">helpers.h:127</a></div></div>
505<div class="ttc" id="normalization__layer_8cl_xhtml_ac0bf948069436897cf005ea45ac73a4f"><div class="ttname"><a href="normalization__layer_8cl.xhtml#ac0bf948069436897cf005ea45ac73a4f">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00026">normalization_layer.cl:26</a></div></div>
506<div class="ttc" id="normalization__layer_8cl_xhtml_a10a9bf1a9fce27664d5f91844ff3f691"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a10a9bf1a9fce27664d5f91844ff3f691">LOAD_OP</a></div><div class="ttdeci">#define LOAD_OP(offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00032">normalization_layer.cl:32</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]507<div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]508<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ad366eaa4cf2d106037d91c30795a5c76"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">arm_compute::test::validation::in</a></div><div class="ttdeci">CLTensor in</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_winograd_8cpp_source.xhtml#l00241">Winograd.cpp:241</a></div></div>
509<div class="ttc" id="normalization__layer_8cl_xhtml_a107d847044e677b01e9bd3d5251b39d9"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a></div><div class="ttdeci">#define SQCVT_SAT(a)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00030">normalization_layer.cl:30</a></div></div>
510<div class="ttc" id="struct_tensor3_d_xhtml"><div class="ttname"><a href="struct_tensor3_d.xhtml">Tensor3D</a></div><div class="ttdoc">Structure to hold 3D tensor information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00159">helpers.h:159</a></div></div>
511<div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00056">helpers.h:56</a></div></div>
512<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_af8a8a0625e7981212a0af48deb9d2a09"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#af8a8a0625e7981212a0af48deb9d2a09">arm_compute::test::validation::out</a></div><div class="ttdeci">CLTensor out</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_winograd_8cpp_source.xhtml#l00242">Winograd.cpp:242</a></div></div>
513<div class="ttc" id="normalization__layer_8cl_xhtml_ac59caf44b60378c02032aab994ad5ac1"><div class="ttname"><a href="normalization__layer_8cl.xhtml#ac59caf44b60378c02032aab994ad5ac1">STORE_OP</a></div><div class="ttdeci">#define STORE_OP(data, offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00033">normalization_layer.cl:33</a></div></div>
514<div class="ttc" id="helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00323">helpers.h:323</a></div></div>
515<div class="ttc" id="normalization__layer_8cl_xhtml_a334426d638df3b4108c366e240cf5504"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a334426d638df3b4108c366e240cf5504">ADD_OP</a></div><div class="ttdeci">#define ADD_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00027">normalization_layer.cl:27</a></div></div>
516<div class="ttc" id="normalization__layer_8cl_xhtml_a5db17889d824975fefb2ce2f4690637f"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a></div><div class="ttdeci">#define DIV_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00028">normalization_layer.cl:28</a></div></div>
517<div class="ttc" id="normalization__layer_8cl_xhtml_a09724da1a9e5ba54260905829ec27db4"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a09724da1a9e5ba54260905829ec27db4">POW_OP</a></div><div class="ttdeci">#define POW_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00029">normalization_layer.cl:29</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]518</div><!-- fragment -->
519</div>
520</div>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]521<a class="anchor" id="a98051a32b8ae24598346306c1b131b2c"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]522<div class="memitem">
523<div class="memproto">
524 <table class="memname">
525 <tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]526 <td class="memname">__kernel void normalization_layer_in_map </td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]527 <td>(</td>
528 <td class="paramtype">__global uchar *&#160;</td>
529 <td class="paramname"><em>input_ptr</em>, </td>
530 </tr>
531 <tr>
532 <td class="paramkey"></td>
533 <td></td>
534 <td class="paramtype">uint&#160;</td>
535 <td class="paramname"><em>input_stride_x</em>, </td>
536 </tr>
537 <tr>
538 <td class="paramkey"></td>
539 <td></td>
540 <td class="paramtype">uint&#160;</td>
541 <td class="paramname"><em>input_step_x</em>, </td>
542 </tr>
543 <tr>
544 <td class="paramkey"></td>
545 <td></td>
546 <td class="paramtype">uint&#160;</td>
547 <td class="paramname"><em>input_stride_y</em>, </td>
548 </tr>
549 <tr>
550 <td class="paramkey"></td>
551 <td></td>
552 <td class="paramtype">uint&#160;</td>
553 <td class="paramname"><em>input_step_y</em>, </td>
554 </tr>
555 <tr>
556 <td class="paramkey"></td>
557 <td></td>
558 <td class="paramtype">uint&#160;</td>
559 <td class="paramname"><em>input_stride_z</em>, </td>
560 </tr>
561 <tr>
562 <td class="paramkey"></td>
563 <td></td>
564 <td class="paramtype">uint&#160;</td>
565 <td class="paramname"><em>input_step_z</em>, </td>
566 </tr>
567 <tr>
568 <td class="paramkey"></td>
569 <td></td>
570 <td class="paramtype">uint&#160;</td>
571 <td class="paramname"><em>input_offset_first_element_in_bytes</em>, </td>
572 </tr>
573 <tr>
574 <td class="paramkey"></td>
575 <td></td>
576 <td class="paramtype">__global uchar *&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]577 <td class="paramname"><em>output_ptr</em>, </td>
578 </tr>
579 <tr>
580 <td class="paramkey"></td>
581 <td></td>
582 <td class="paramtype">uint&#160;</td>
583 <td class="paramname"><em>output_stride_x</em>, </td>
584 </tr>
585 <tr>
586 <td class="paramkey"></td>
587 <td></td>
588 <td class="paramtype">uint&#160;</td>
589 <td class="paramname"><em>output_step_x</em>, </td>
590 </tr>
591 <tr>
592 <td class="paramkey"></td>
593 <td></td>
594 <td class="paramtype">uint&#160;</td>
595 <td class="paramname"><em>output_stride_y</em>, </td>
596 </tr>
597 <tr>
598 <td class="paramkey"></td>
599 <td></td>
600 <td class="paramtype">uint&#160;</td>
601 <td class="paramname"><em>output_step_y</em>, </td>
602 </tr>
603 <tr>
604 <td class="paramkey"></td>
605 <td></td>
606 <td class="paramtype">uint&#160;</td>
607 <td class="paramname"><em>output_stride_z</em>, </td>
608 </tr>
609 <tr>
610 <td class="paramkey"></td>
611 <td></td>
612 <td class="paramtype">uint&#160;</td>
613 <td class="paramname"><em>output_step_z</em>, </td>
614 </tr>
615 <tr>
616 <td class="paramkey"></td>
617 <td></td>
618 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]619 <td class="paramname"><em>output_offset_first_element_in_bytes</em>&#160;</td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]620 </tr>
621 <tr>
622 <td></td>
623 <td>)</td>
624 <td></td><td></td>
625 </tr>
626 </table>
627</div><div class="memdoc">
628
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]629<p>Apply in-map normalization. </p>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]630<dl class="section note"><dt>Note</dt><dd>Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short </dd>
631<dd>
632<a class="el" href="struct_vector.xhtml" title="Structure to hold Vector information. ">Vector</a> size should be given as a preprocessor argument using -DVEC_SIZE=size, e.g. -DVEC_SIZE=16 </dd>
633<dd>
634The radius should be given as a preprocessor argument using -DRADIUS=size. e.g. -DRADIUS=5 </dd>
635<dd>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]636Scaling coefficient (= alpha/norm_size), beta and kappa need to be passed at compile time using -DCOEFF, -DALPHA and -DKAPPA</dd></dl>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]637<dl class="params"><dt>Parameters</dt><dd>
638 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]639 <tr><td class="paramdir">[in]</td><td class="paramname">input_ptr</td><td>Pointer to the first source tensor. Supported data types: F16/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]640 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_x</td><td>Stride of the first source tensor in X dimension (in bytes) </td></tr>
641 <tr><td class="paramdir">[in]</td><td class="paramname">input_step_x</td><td>input_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
642 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_y</td><td>Stride of the first source tensor in Y dimension (in bytes) </td></tr>
643 <tr><td class="paramdir">[in]</td><td class="paramname">input_step_y</td><td>input_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
644 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_z</td><td>Stride of the first source tensor in Z dimension (in bytes) </td></tr>
645 <tr><td class="paramdir">[in]</td><td class="paramname">input_step_z</td><td>input_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
646 <tr><td class="paramdir">[in]</td><td class="paramname">input_offset_first_element_in_bytes</td><td>The offset of the first element in the first source tensor </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]647 <tr><td class="paramdir">[out]</td><td class="paramname">output_ptr</td><td>Pointer to the destination tensor. Supported data types: same as <code>input_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]648 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_x</td><td>Stride of the destination tensor in X dimension (in bytes) </td></tr>
649 <tr><td class="paramdir">[in]</td><td class="paramname">output_step_x</td><td>output_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
650 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_y</td><td>Stride of the first destination tensor in Y dimension (in bytes) </td></tr>
651 <tr><td class="paramdir">[in]</td><td class="paramname">output_step_y</td><td>output_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
652 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_z</td><td>Stride of the first source tensor in Z dimension (in bytes) </td></tr>
653 <tr><td class="paramdir">[in]</td><td class="paramname">output_step_z</td><td>output_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
654 <tr><td class="paramdir">[in]</td><td class="paramname">output_offset_first_element_in_bytes</td><td>The offset of the first element in the destination tensor </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]655 </table>
656 </dd>
657</dl>
658
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]659<p>Definition at line <a class="el" href="normalization__layer_8cl_source.xhtml#l00119">119</a> of file <a class="el" href="normalization__layer_8cl_source.xhtml">normalization_layer.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]660
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]661<p>References <a class="el" href="normalization__layer_8cl_source.xhtml#l00027">ADD_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00127">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00028">DIV_OP</a>, <a class="el" href="_c_l_2_winograd_8cpp_source.xhtml#l00241">arm_compute::test::validation::in</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00032">LOAD_OP</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00026">MUL_OP</a>, <a class="el" href="_c_l_2_winograd_8cpp_source.xhtml#l00242">arm_compute::test::validation::out</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00029">POW_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00161">Tensor3D::ptr</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00030">SQCVT_SAT</a>, <a class="el" href="normalization__layer_8cl_source.xhtml#l00033">STORE_OP</a>, <a class="el" href="helpers_8h_source.xhtml#l00323">tensor3D_offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00056">VEC_DATA_TYPE</a>.</p>
662<div class="fragment"><div class="line"><a name="l00121"></a><span class="lineno"> 121</span>&#160;{</div><div class="line"><a name="l00122"></a><span class="lineno"> 122</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">in</a> = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(input);</div><div class="line"><a name="l00123"></a><span class="lineno"> 123</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#af8a8a0625e7981212a0af48deb9d2a09">out</a> = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(output);</div><div class="line"><a name="l00124"></a><span class="lineno"> 124</span>&#160;</div><div class="line"><a name="l00125"></a><span class="lineno"> 125</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, VEC_SIZE)</div><div class="line"><a name="l00126"></a><span class="lineno"> 126</span>&#160; acc = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, VEC_SIZE))0;</div><div class="line"><a name="l00127"></a><span class="lineno"> 127</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00128"></a><span class="lineno"> 128</span>&#160; coeff_v = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE))<a class="code" href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a>(COEFF);</div><div class="line"><a name="l00129"></a><span class="lineno"> 129</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00130"></a><span class="lineno"> 130</span>&#160; beta_v = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE))<a class="code" href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a>(BETA);</div><div class="line"><a name="l00131"></a><span class="lineno"> 131</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00132"></a><span class="lineno"> 132</span>&#160; kappa_v = (<a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE))<a class="code" href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a>(KAPPA);</div><div class="line"><a name="l00133"></a><span class="lineno"> 133</span>&#160;</div><div class="line"><a name="l00134"></a><span class="lineno"> 134</span>&#160; const <span class="keywordtype">int</span> current_col = get_global_id(0) &lt;&lt; 2;</div><div class="line"><a name="l00135"></a><span class="lineno"> 135</span>&#160; const <span class="keywordtype">int</span> left_pos = max(-(<span class="keywordtype">int</span>)RADIUS, -3 - current_col);</div><div class="line"><a name="l00136"></a><span class="lineno"> 136</span>&#160; const <span class="keywordtype">int</span> right_pos = min((<span class="keywordtype">int</span>)RADIUS, (<span class="keywordtype">int</span>)((get_global_size(0) &lt;&lt; 2) + 3 - 1 - current_col));</div><div class="line"><a name="l00137"></a><span class="lineno"> 137</span>&#160;</div><div class="line"><a name="l00138"></a><span class="lineno"> 138</span>&#160;<span class="preprocessor">#if defined(IN_MAP_2D)</span></div><div class="line"><a name="l00139"></a><span class="lineno"> 139</span>&#160; <span class="keyword">const</span> <span class="keywordtype">int</span> current_row = get_global_id(1);</div><div class="line"><a name="l00140"></a><span class="lineno"> 140</span>&#160; <span class="keyword">const</span> <span class="keywordtype">int</span> first_row = max(-(<span class="keywordtype">int</span>)RADIUS, -current_row);</div><div class="line"><a name="l00141"></a><span class="lineno"> 141</span>&#160; <span class="keyword">const</span> <span class="keywordtype">int</span> last_row = min((<span class="keywordtype">int</span>)RADIUS, (<span class="keywordtype">int</span>)get_global_size(1) - 1 - current_row);</div><div class="line"><a name="l00142"></a><span class="lineno"> 142</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(IN_MAP_2D) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00143"></a><span class="lineno"> 143</span>&#160;</div><div class="line"><a name="l00144"></a><span class="lineno"> 144</span>&#160;<span class="preprocessor">#if defined(IN_MAP_2D)</span></div><div class="line"><a name="l00145"></a><span class="lineno"> 145</span>&#160; <span class="keywordflow">for</span>(<span class="keywordtype">int</span> j = first_row; j &lt;= last_row; ++j)</div><div class="line"><a name="l00146"></a><span class="lineno"> 146</span>&#160; {</div><div class="line"><a name="l00147"></a><span class="lineno"> 147</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(IN_MAP_2D) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00148"></a><span class="lineno"> 148</span>&#160; <span class="keywordflow">for</span>(<span class="keywordtype">int</span> i = left_pos; i &lt;= right_pos; ++i)</div><div class="line"><a name="l00149"></a><span class="lineno"> 149</span>&#160; {</div><div class="line"><a name="l00150"></a><span class="lineno"> 150</span>&#160;<span class="preprocessor">#if defined(IN_MAP_2D)</span></div><div class="line"><a name="l00151"></a><span class="lineno"> 151</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span>&#160; values = <a class="code" href="normalization__layer_8cl.xhtml#a10a9bf1a9fce27664d5f91844ff3f691">LOAD_OP</a>(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;in, i, j, 0));</div><div class="line"><a name="l00153"></a><span class="lineno"> 153</span>&#160;<span class="preprocessor">#else </span><span class="comment">/* defined(IN_MAP_2D) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00154"></a><span class="lineno"> 154</span>&#160; <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span>&#160; values = LOAD_OP(0, (__global DATA_TYPE *)<a class="code" href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;in, i, 0, 0));</div><div class="line"><a name="l00156"></a><span class="lineno"> 156</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(IN_MAP_2D) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00157"></a><span class="lineno"> 157</span>&#160; acc = <a class="code" href="normalization__layer_8cl.xhtml#a334426d638df3b4108c366e240cf5504">ADD_OP</a>(acc, <a class="code" href="normalization__layer_8cl.xhtml#ac0bf948069436897cf005ea45ac73a4f">MUL_OP</a>(values, values));</div><div class="line"><a name="l00158"></a><span class="lineno"> 158</span>&#160; }</div><div class="line"><a name="l00159"></a><span class="lineno"> 159</span>&#160;<span class="preprocessor">#if defined(IN_MAP_2D)</span></div><div class="line"><a name="l00160"></a><span class="lineno"> 160</span>&#160; }</div><div class="line"><a name="l00161"></a><span class="lineno"> 161</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(IN_MAP_2D) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00162"></a><span class="lineno"> 162</span>&#160;</div><div class="line"><a name="l00163"></a><span class="lineno"> 163</span>&#160; acc = <a class="code" href="normalization__layer_8cl.xhtml#a334426d638df3b4108c366e240cf5504">ADD_OP</a>(<a class="code" href="normalization__layer_8cl.xhtml#ac0bf948069436897cf005ea45ac73a4f">MUL_OP</a>(acc, coeff_v), kappa_v);</div><div class="line"><a name="l00164"></a><span class="lineno"> 164</span>&#160; <span class="keyword">const</span> <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00165"></a><span class="lineno"> 165</span>&#160; normalized = <a class="code" href="normalization__layer_8cl.xhtml#a09724da1a9e5ba54260905829ec27db4">POW_OP</a>(acc, beta_v);</div><div class="line"><a name="l00166"></a><span class="lineno"> 166</span>&#160; const <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE, VEC_SIZE)</div><div class="line"><a name="l00167"></a><span class="lineno"> 167</span>&#160; normalized_pixel = <a class="code" href="normalization__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a>(LOAD_OP(0, (__global DATA_TYPE *)in.ptr), normalized);</div><div class="line"><a name="l00168"></a><span class="lineno"> 168</span>&#160;</div><div class="line"><a name="l00169"></a><span class="lineno"> 169</span>&#160; <a class="code" href="normalization__layer_8cl.xhtml#ac59caf44b60378c02032aab994ad5ac1">STORE_OP</a>(normalized_pixel, 0, (__global DATA_TYPE *)out.ptr);</div><div class="line"><a name="l00170"></a><span class="lineno"> 170</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00127">helpers.h:127</a></div></div>
663<div class="ttc" id="normalization__layer_8cl_xhtml_ac0bf948069436897cf005ea45ac73a4f"><div class="ttname"><a href="normalization__layer_8cl.xhtml#ac0bf948069436897cf005ea45ac73a4f">MUL_OP</a></div><div class="ttdeci">#define MUL_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00026">normalization_layer.cl:26</a></div></div>
664<div class="ttc" id="normalization__layer_8cl_xhtml_a10a9bf1a9fce27664d5f91844ff3f691"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a10a9bf1a9fce27664d5f91844ff3f691">LOAD_OP</a></div><div class="ttdeci">#define LOAD_OP(offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00032">normalization_layer.cl:32</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]665<div class="ttc" id="convolution3x3_8cl_xhtml_afb8c72ce35c4a1f4a2588d6573e54aa1"><div class="ttname"><a href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a></div><div class="ttdeci">#define DATA_TYPE</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00027">convolution3x3.cl:27</a></div></div>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]666<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ad366eaa4cf2d106037d91c30795a5c76"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ad366eaa4cf2d106037d91c30795a5c76">arm_compute::test::validation::in</a></div><div class="ttdeci">CLTensor in</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_winograd_8cpp_source.xhtml#l00241">Winograd.cpp:241</a></div></div>
667<div class="ttc" id="normalization__layer_8cl_xhtml_a107d847044e677b01e9bd3d5251b39d9"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a107d847044e677b01e9bd3d5251b39d9">SQCVT_SAT</a></div><div class="ttdeci">#define SQCVT_SAT(a)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00030">normalization_layer.cl:30</a></div></div>
668<div class="ttc" id="struct_tensor3_d_xhtml"><div class="ttname"><a href="struct_tensor3_d.xhtml">Tensor3D</a></div><div class="ttdoc">Structure to hold 3D tensor information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00159">helpers.h:159</a></div></div>
669<div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00056">helpers.h:56</a></div></div>
670<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_af8a8a0625e7981212a0af48deb9d2a09"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#af8a8a0625e7981212a0af48deb9d2a09">arm_compute::test::validation::out</a></div><div class="ttdeci">CLTensor out</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_winograd_8cpp_source.xhtml#l00242">Winograd.cpp:242</a></div></div>
671<div class="ttc" id="normalization__layer_8cl_xhtml_ac59caf44b60378c02032aab994ad5ac1"><div class="ttname"><a href="normalization__layer_8cl.xhtml#ac59caf44b60378c02032aab994ad5ac1">STORE_OP</a></div><div class="ttdeci">#define STORE_OP(data, offset, ptr)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00033">normalization_layer.cl:33</a></div></div>
672<div class="ttc" id="helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00323">helpers.h:323</a></div></div>
673<div class="ttc" id="normalization__layer_8cl_xhtml_a334426d638df3b4108c366e240cf5504"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a334426d638df3b4108c366e240cf5504">ADD_OP</a></div><div class="ttdeci">#define ADD_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00027">normalization_layer.cl:27</a></div></div>
674<div class="ttc" id="normalization__layer_8cl_xhtml_a5db17889d824975fefb2ce2f4690637f"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a5db17889d824975fefb2ce2f4690637f">DIV_OP</a></div><div class="ttdeci">#define DIV_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00028">normalization_layer.cl:28</a></div></div>
675<div class="ttc" id="normalization__layer_8cl_xhtml_a09724da1a9e5ba54260905829ec27db4"><div class="ttname"><a href="normalization__layer_8cl.xhtml#a09724da1a9e5ba54260905829ec27db4">POW_OP</a></div><div class="ttdeci">#define POW_OP(x, y)</div><div class="ttdef"><b>Definition:</b> <a href="normalization__layer_8cl_source.xhtml#l00029">normalization_layer.cl:29</a></div></div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]676</div><!-- fragment -->
677</div>
678</div>
679</div><!-- contents -->
680</div><!-- doc-content -->
681<!-- start footer part -->
682<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
683 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]684 <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="normalization__layer_8cl.xhtml">normalization_layer.cl</a></li>
Jenkins52ba29e2018-08-29 15:32:11 +0000[diff] [blame^]685 <li class="footer">Generated on Wed Aug 29 2018 15:31:38 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]686 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]687 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]688 </ul>
689</div>
690</body>
691</html>