blob: 35f11919cd3e0856bff72cdfbfc29adc128321c3 [file] [log] [blame]
Jenkins514be652019-02-28 12:25:18 +00001<!-- HTML header for doxygen 1.8.15-->
2<!-- Remember to use version doxygen 1.8.15 +-->
Kaizen8938bd32017-09-28 14:38:23 +01003<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
4<html xmlns="http://www.w3.org/1999/xhtml">
5<head>
6<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
7<meta http-equiv="X-UA-Compatible" content="IE=9"/>
Jenkins514be652019-02-28 12:25:18 +00008<meta name="generator" content="Doxygen 1.8.15"/>
Kaizen8938bd32017-09-28 14:38:23 +01009<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
10<title>Compute Library: src/core/CL/cl_kernels/l2_normalize.cl File Reference</title>
11<link href="tabs.css" rel="stylesheet" type="text/css"/>
12<script type="text/javascript" src="jquery.js"></script>
13<script type="text/javascript" src="dynsections.js"></script>
14<link href="navtree.css" rel="stylesheet" type="text/css"/>
15<script type="text/javascript" src="resize.js"></script>
Anthony Barbier8140e1e2017-12-14 23:48:46 +000016<script type="text/javascript" src="navtreedata.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010017<script type="text/javascript" src="navtree.js"></script>
18<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +000019/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Kaizen8938bd32017-09-28 14:38:23 +010020 $(document).ready(initResizable);
Jenkins514be652019-02-28 12:25:18 +000021/* @license-end */</script>
Kaizen8938bd32017-09-28 14:38:23 +010022<link href="search/search.css" rel="stylesheet" type="text/css"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +000023<script type="text/javascript" src="search/searchdata.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010024<script type="text/javascript" src="search/search.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010025<script type="text/x-mathjax-config">
26 MathJax.Hub.Config({
27 extensions: ["tex2jax.js"],
28 jax: ["input/TeX","output/HTML-CSS"],
29});
Jenkins514be652019-02-28 12:25:18 +000030</script><script type="text/javascript" async="async" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010031<link href="doxygen.css" rel="stylesheet" type="text/css" />
Jenkins514be652019-02-28 12:25:18 +000032<link href="stylesheet.css" rel="stylesheet" type="text/css"/>
Kaizen8938bd32017-09-28 14:38:23 +010033</head>
34<body>
35<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
36<div id="titlearea">
37<table cellspacing="0" cellpadding="0">
38 <tbody>
39 <tr style="height: 56px;">
Jenkins514be652019-02-28 12:25:18 +000040 <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"/>
Kaizen8938bd32017-09-28 14:38:23 +010041 <td style="padding-left: 0.5em;">
Jenkins514be652019-02-28 12:25:18 +000042 <div id="projectname">
43 &#160;<span id="projectnumber">19.02</span>
Kaizen8938bd32017-09-28 14:38:23 +010044 </div>
45 </td>
46 </tr>
47 </tbody>
48</table>
49</div>
50<!-- end header part -->
Jenkins514be652019-02-28 12:25:18 +000051<!-- Generated by Doxygen 1.8.15 -->
Kaizen8938bd32017-09-28 14:38:23 +010052<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +000053/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Kaizen8938bd32017-09-28 14:38:23 +010054var searchBox = new SearchBox("searchBox", "search",false,'Search');
Jenkins514be652019-02-28 12:25:18 +000055/* @license-end */
Kaizen8938bd32017-09-28 14:38:23 +010056</script>
Jenkinsb9abeae2018-11-22 11:58:08 +000057<script type="text/javascript" src="menudata.js"></script>
58<script type="text/javascript" src="menu.js"></script>
59<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +000060/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Jenkinsb9abeae2018-11-22 11:58:08 +000061$(function() {
62 initMenu('',true,false,'search.php','Search');
63 $(document).ready(function() { init_search(); });
64});
Jenkins514be652019-02-28 12:25:18 +000065/* @license-end */</script>
Jenkinsb9abeae2018-11-22 11:58:08 +000066<div id="main-nav"></div>
Kaizen8938bd32017-09-28 14:38:23 +010067</div><!-- top -->
68<div id="side-nav" class="ui-resizable side-nav-resizable">
69 <div id="nav-tree">
70 <div id="nav-tree-contents">
71 <div id="nav-sync" class="sync"></div>
72 </div>
73 </div>
74 <div id="splitbar" style="-moz-user-select:none;"
75 class="ui-resizable-handle">
76 </div>
77</div>
78<script type="text/javascript">
Jenkins514be652019-02-28 12:25:18 +000079/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
Kaizen8938bd32017-09-28 14:38:23 +010080$(document).ready(function(){initNavTree('l2__normalize_8cl.xhtml','');});
Jenkins514be652019-02-28 12:25:18 +000081/* @license-end */
Kaizen8938bd32017-09-28 14:38:23 +010082</script>
83<div id="doc-content">
84<!-- window showing the filter options -->
85<div id="MSearchSelectWindow"
86 onmouseover="return searchBox.OnSearchSelectShow()"
87 onmouseout="return searchBox.OnSearchSelectHide()"
88 onkeydown="return searchBox.OnSearchSelectKey(event)">
Anthony Barbier8140e1e2017-12-14 23:48:46 +000089</div>
Kaizen8938bd32017-09-28 14:38:23 +010090
91<!-- iframe showing the search results (closed by default) -->
92<div id="MSearchResultsWindow">
93<iframe src="javascript:void(0)" frameborder="0"
94 name="MSearchResults" id="MSearchResults">
95</iframe>
96</div>
97
98<div class="header">
99 <div class="summary">
100<a href="#func-members">Functions</a> </div>
101 <div class="headertitle">
102<div class="title">l2_normalize.cl File Reference</div> </div>
103</div><!--header-->
104<div class="contents">
Jenkinsb9abeae2018-11-22 11:58:08 +0000105<div class="textblock"><code>#include &quot;<a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
Kaizen8938bd32017-09-28 14:38:23 +0100106</div>
107<p><a href="l2__normalize_8cl_source.xhtml">Go to the source code of this file.</a></p>
108<table class="memberdecls">
109<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
110Functions</h2></td></tr>
Jenkinsb9abeae2018-11-22 11:58:08 +0000111<tr class="memitem:aecb8206d06294af64d2cf722399df448"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="l2__normalize_8cl.xhtml#aecb8206d06294af64d2cf722399df448">l2_normalize_x</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_offset_first_element_in_bytes, __global uchar *sum_ptr, uint sum_stride_x, uint sum_step_x, uint sum_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_offset_first_element_in_bytes, <a class="el" href="convolution9x9_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> <a class="el" href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a>)</td></tr>
112<tr class="memdesc:aecb8206d06294af64d2cf722399df448"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel performs l2 normalization on x-axis. <a href="#aecb8206d06294af64d2cf722399df448">More...</a><br /></td></tr>
113<tr class="separator:aecb8206d06294af64d2cf722399df448"><td class="memSeparator" colspan="2">&#160;</td></tr>
114<tr class="memitem:aa860f55bdbc8de4e8e5850a05147b2be"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="l2__normalize_8cl.xhtml#aa860f55bdbc8de4e8e5850a05147b2be">l2_normalize_y</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, 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_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_offset_first_element_in_bytes, <a class="el" href="convolution9x9_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> <a class="el" href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a>)</td></tr>
115<tr class="memdesc:aa860f55bdbc8de4e8e5850a05147b2be"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel performs l2 normalization on y-axis. <a href="#aa860f55bdbc8de4e8e5850a05147b2be">More...</a><br /></td></tr>
116<tr class="separator:aa860f55bdbc8de4e8e5850a05147b2be"><td class="memSeparator" colspan="2">&#160;</td></tr>
117<tr class="memitem:a809db37fce4c2e573b09b0af0d3b95ee"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="l2__normalize_8cl.xhtml#a809db37fce4c2e573b09b0af0d3b95ee">l2_normalize_z</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, <a class="el" href="convolution9x9_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> <a class="el" href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a>)</td></tr>
118<tr class="memdesc:a809db37fce4c2e573b09b0af0d3b95ee"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel performs l2 normalization on z-axis. <a href="#a809db37fce4c2e573b09b0af0d3b95ee">More...</a><br /></td></tr>
119<tr class="separator:a809db37fce4c2e573b09b0af0d3b95ee"><td class="memSeparator" colspan="2">&#160;</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100120</table>
121<h2 class="groupheader">Function Documentation</h2>
Jenkinsb9abeae2018-11-22 11:58:08 +0000122<a id="aecb8206d06294af64d2cf722399df448"></a>
123<h2 class="memtitle"><span class="permalink"><a href="#aecb8206d06294af64d2cf722399df448">&#9670;&nbsp;</a></span>l2_normalize_x()</h2>
124
Kaizen8938bd32017-09-28 14:38:23 +0100125<div class="memitem">
126<div class="memproto">
127 <table class="memname">
128 <tr>
Jenkinsb9abeae2018-11-22 11:58:08 +0000129 <td class="memname">__kernel void l2_normalize_x </td>
Kaizen8938bd32017-09-28 14:38:23 +0100130 <td>(</td>
131 <td class="paramtype">__global uchar *&#160;</td>
132 <td class="paramname"><em>src_ptr</em>, </td>
133 </tr>
134 <tr>
135 <td class="paramkey"></td>
136 <td></td>
137 <td class="paramtype">uint&#160;</td>
138 <td class="paramname"><em>src_stride_x</em>, </td>
139 </tr>
140 <tr>
141 <td class="paramkey"></td>
142 <td></td>
143 <td class="paramtype">uint&#160;</td>
144 <td class="paramname"><em>src_step_x</em>, </td>
145 </tr>
146 <tr>
147 <td class="paramkey"></td>
148 <td></td>
149 <td class="paramtype">uint&#160;</td>
150 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
151 </tr>
152 <tr>
153 <td class="paramkey"></td>
154 <td></td>
155 <td class="paramtype">__global uchar *&#160;</td>
156 <td class="paramname"><em>sum_ptr</em>, </td>
157 </tr>
158 <tr>
159 <td class="paramkey"></td>
160 <td></td>
161 <td class="paramtype">uint&#160;</td>
162 <td class="paramname"><em>sum_stride_x</em>, </td>
163 </tr>
164 <tr>
165 <td class="paramkey"></td>
166 <td></td>
167 <td class="paramtype">uint&#160;</td>
168 <td class="paramname"><em>sum_step_x</em>, </td>
169 </tr>
170 <tr>
171 <td class="paramkey"></td>
172 <td></td>
173 <td class="paramtype">uint&#160;</td>
174 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
175 </tr>
176 <tr>
177 <td class="paramkey"></td>
178 <td></td>
179 <td class="paramtype">__global uchar *&#160;</td>
180 <td class="paramname"><em>dst_ptr</em>, </td>
181 </tr>
182 <tr>
183 <td class="paramkey"></td>
184 <td></td>
185 <td class="paramtype">uint&#160;</td>
186 <td class="paramname"><em>dst_stride_x</em>, </td>
187 </tr>
188 <tr>
189 <td class="paramkey"></td>
190 <td></td>
191 <td class="paramtype">uint&#160;</td>
192 <td class="paramname"><em>dst_step_x</em>, </td>
193 </tr>
194 <tr>
195 <td class="paramkey"></td>
196 <td></td>
197 <td class="paramtype">uint&#160;</td>
198 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
199 </tr>
200 <tr>
201 <td class="paramkey"></td>
202 <td></td>
203 <td class="paramtype"><a class="el" href="convolution9x9_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>&#160;</td>
204 <td class="paramname"><em>epsilon</em>&#160;</td>
205 </tr>
206 <tr>
207 <td></td>
208 <td>)</td>
209 <td></td><td></td>
210 </tr>
211 </table>
212</div><div class="memdoc">
213
Jenkinsb9abeae2018-11-22 11:58:08 +0000214<p>This kernel performs l2 normalization on x-axis. </p>
Kaizen8938bd32017-09-28 14:38:23 +0100215<dl class="section note"><dt>Note</dt><dd>The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float </dd>
216<dd>
217The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32</dd></dl>
218<dl class="params"><dt>Parameters</dt><dd>
219 <table class="params">
Jenkins52ba29e2018-08-29 15:32:11 +0000220 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: F16/F32 </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100221 <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>
222 <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>
223 <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>
Jenkins52ba29e2018-08-29 15:32:11 +0000224 <tr><td class="paramdir">[in]</td><td class="paramname">sum_ptr</td><td>Pointer to the source tensor. Supported data types: F16/F32 </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100225 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_x</td><td>Stride of the source tensor in X dimension (in bytes) </td></tr>
226 <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>
227 <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 source tensor </td></tr>
228 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: same as <code>src_ptr</code> </td></tr>
229 <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>
230 <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>
231 <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>
232 <tr><td class="paramdir">[in]</td><td class="paramname">epsilon</td><td>Epsilon value </td></tr>
233 </table>
234 </dd>
235</dl>
236
Jenkinsb9abeae2018-11-22 11:58:08 +0000237<p class="definition">Definition at line <a class="el" href="l2__normalize_8cl_source.xhtml#l00045">45</a> of file <a class="el" href="l2__normalize_8cl_source.xhtml">l2_normalize.cl</a>.</p>
Jenkins514be652019-02-28 12:25:18 +0000238<div class="fragment"><div class="line"><a name="l00050"></a><span class="lineno"> 50</span>&#160;{</div><div class="line"><a name="l00051"></a><span class="lineno"> 51</span>&#160; <a class="code" href="struct_vector.xhtml">Vector</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a>);</div><div class="line"><a name="l00052"></a><span class="lineno"> 52</span>&#160; <a class="code" href="struct_vector.xhtml">Vector</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#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a>(<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>);</div><div class="line"><a name="l00053"></a><span class="lineno"> 53</span>&#160; <a class="code" href="struct_vector.xhtml">Vector</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#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00054"></a><span class="lineno"> 54</span>&#160;</div><div class="line"><a name="l00055"></a><span class="lineno"> 55</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00056"></a><span class="lineno"> 56</span>&#160; in = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a>.ptr);</div><div class="line"><a name="l00057"></a><span class="lineno"> 57</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>&#160; normalize_value = (<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))rsqrt(fmax(((__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>.ptr)[0], <a class="code" href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a>));</div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>&#160;</div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span>&#160; vstore16(in * normalize_value, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>.ptr);</div><div class="line"><a name="l00061"></a><span class="lineno"> 61</span>&#160;}</div><div class="ttc" id="struct_vector_xhtml"><div class="ttname"><a href="struct_vector.xhtml">Vector</a></div><div class="ttdoc">Structure to hold Vector information.</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00143">helpers.h:143</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100239<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>
Jenkins514be652019-02-28 12:25:18 +0000240<div class="ttc" id="reduction__operation_8cl_xhtml_ab0df00f5333da51860deb93deb44a782"><div class="ttname"><a href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a></div><div class="ttdeci">DATA_TYPE sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate sum of a vector.</div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00052">reduction_operation.cl:52</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000241<div class="ttc" id="_asymm_helpers_8cpp_xhtml_a552dc3787d7ea1675f3e4e8993501d58"><div class="ttname"><a href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a></div><div class="ttdeci">constexpr float epsilon</div><div class="ttdef"><b>Definition:</b> <a href="_asymm_helpers_8cpp_source.xhtml#l00033">AsymmHelpers.cpp:33</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000242<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a527bfdf5eeb306f1cf01c4a8e29f38e0"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a527bfdf5eeb306f1cf01c4a8e29f38e0">CONVERT_TO_VECTOR_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00107">helpers.h:107</a></div></div>
243<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>
244<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ae9d2dc29c2789c253406f9b304cc75a8"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
245<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#l00057">helpers.h:57</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000246</div><!-- fragment -->
Jenkins514be652019-02-28 12:25:18 +0000247<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00107">CONVERT_TO_VECTOR_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="_asymm_helpers_8cpp_source.xhtml#l00033">epsilon</a>, <a class="el" href="_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00057">VEC_DATA_TYPE</a>.</p>
248
Jenkinsb9abeae2018-11-22 11:58:08 +0000249</div>
250</div>
251<a id="aa860f55bdbc8de4e8e5850a05147b2be"></a>
252<h2 class="memtitle"><span class="permalink"><a href="#aa860f55bdbc8de4e8e5850a05147b2be">&#9670;&nbsp;</a></span>l2_normalize_y()</h2>
253
254<div class="memitem">
255<div class="memproto">
256 <table class="memname">
257 <tr>
258 <td class="memname">__kernel void l2_normalize_y </td>
259 <td>(</td>
260 <td class="paramtype">__global uchar *&#160;</td>
261 <td class="paramname"><em>src_ptr</em>, </td>
262 </tr>
263 <tr>
264 <td class="paramkey"></td>
265 <td></td>
266 <td class="paramtype">uint&#160;</td>
267 <td class="paramname"><em>src_stride_x</em>, </td>
268 </tr>
269 <tr>
270 <td class="paramkey"></td>
271 <td></td>
272 <td class="paramtype">uint&#160;</td>
273 <td class="paramname"><em>src_step_x</em>, </td>
274 </tr>
275 <tr>
276 <td class="paramkey"></td>
277 <td></td>
278 <td class="paramtype">uint&#160;</td>
279 <td class="paramname"><em>src_stride_y</em>, </td>
280 </tr>
281 <tr>
282 <td class="paramkey"></td>
283 <td></td>
284 <td class="paramtype">uint&#160;</td>
285 <td class="paramname"><em>src_step_y</em>, </td>
286 </tr>
287 <tr>
288 <td class="paramkey"></td>
289 <td></td>
290 <td class="paramtype">uint&#160;</td>
291 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
292 </tr>
293 <tr>
294 <td class="paramkey"></td>
295 <td></td>
296 <td class="paramtype">__global uchar *&#160;</td>
297 <td class="paramname"><em>sum_ptr</em>, </td>
298 </tr>
299 <tr>
300 <td class="paramkey"></td>
301 <td></td>
302 <td class="paramtype">uint&#160;</td>
303 <td class="paramname"><em>sum_stride_x</em>, </td>
304 </tr>
305 <tr>
306 <td class="paramkey"></td>
307 <td></td>
308 <td class="paramtype">uint&#160;</td>
309 <td class="paramname"><em>sum_step_x</em>, </td>
310 </tr>
311 <tr>
312 <td class="paramkey"></td>
313 <td></td>
314 <td class="paramtype">uint&#160;</td>
315 <td class="paramname"><em>sum_stride_y</em>, </td>
316 </tr>
317 <tr>
318 <td class="paramkey"></td>
319 <td></td>
320 <td class="paramtype">uint&#160;</td>
321 <td class="paramname"><em>sum_step_y</em>, </td>
322 </tr>
323 <tr>
324 <td class="paramkey"></td>
325 <td></td>
326 <td class="paramtype">uint&#160;</td>
327 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
328 </tr>
329 <tr>
330 <td class="paramkey"></td>
331 <td></td>
332 <td class="paramtype">__global uchar *&#160;</td>
333 <td class="paramname"><em>dst_ptr</em>, </td>
334 </tr>
335 <tr>
336 <td class="paramkey"></td>
337 <td></td>
338 <td class="paramtype">uint&#160;</td>
339 <td class="paramname"><em>dst_stride_x</em>, </td>
340 </tr>
341 <tr>
342 <td class="paramkey"></td>
343 <td></td>
344 <td class="paramtype">uint&#160;</td>
345 <td class="paramname"><em>dst_step_x</em>, </td>
346 </tr>
347 <tr>
348 <td class="paramkey"></td>
349 <td></td>
350 <td class="paramtype">uint&#160;</td>
351 <td class="paramname"><em>dst_stride_y</em>, </td>
352 </tr>
353 <tr>
354 <td class="paramkey"></td>
355 <td></td>
356 <td class="paramtype">uint&#160;</td>
357 <td class="paramname"><em>dst_step_y</em>, </td>
358 </tr>
359 <tr>
360 <td class="paramkey"></td>
361 <td></td>
362 <td class="paramtype">uint&#160;</td>
363 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
364 </tr>
365 <tr>
366 <td class="paramkey"></td>
367 <td></td>
368 <td class="paramtype"><a class="el" href="convolution9x9_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>&#160;</td>
369 <td class="paramname"><em>epsilon</em>&#160;</td>
370 </tr>
371 <tr>
372 <td></td>
373 <td>)</td>
374 <td></td><td></td>
375 </tr>
376 </table>
377</div><div class="memdoc">
378
379<p>This kernel performs l2 normalization on y-axis. </p>
380<dl class="section note"><dt>Note</dt><dd>The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float </dd>
381<dd>
382The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32</dd></dl>
383<dl class="params"><dt>Parameters</dt><dd>
384 <table class="params">
385 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: F16/F32 </td></tr>
386 <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>
387 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_x</td><td>src_stride_x * number of elements along Y processed per workitem(in bytes) </td></tr>
388 <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>
389 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_y</td><td>src_stride_y * number of elements along X processed per workitem(in bytes) </td></tr>
390 <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>
391 <tr><td class="paramdir">[in]</td><td class="paramname">sum_ptr</td><td>Pointer to the source tensor. Supported data types: F16/F32 </td></tr>
392 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_x</td><td>Stride of the source tensor in X dimension (in bytes) </td></tr>
393 <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>
394 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_y</td><td>Stride of the source tensor in Y dimension (in bytes) </td></tr>
395 <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>
396 <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 source tensor </td></tr>
397 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: same as <code>src_ptr</code> </td></tr>
398 <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>
399 <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>
400 <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>
401 <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>
402 <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>
403 <tr><td class="paramdir">[in]</td><td class="paramname">epsilon</td><td>Epsilon value </td></tr>
404 </table>
405 </dd>
406</dl>
407
408<p class="definition">Definition at line <a class="el" href="l2__normalize_8cl_source.xhtml#l00088">88</a> of file <a class="el" href="l2__normalize_8cl_source.xhtml">l2_normalize.cl</a>.</p>
Jenkins514be652019-02-28 12:25:18 +0000409<div class="fragment"><div class="line"><a name="l00093"></a><span class="lineno"> 93</span>&#160;{</div><div class="line"><a name="l00094"></a><span class="lineno"> 94</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a>);</div><div class="line"><a name="l00095"></a><span class="lineno"> 95</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>);</div><div class="line"><a name="l00096"></a><span class="lineno"> 96</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00097"></a><span class="lineno"> 97</span>&#160;</div><div class="line"><a name="l00098"></a><span class="lineno"> 98</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00099"></a><span class="lineno"> 99</span>&#160; in = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a>.ptr);</div><div class="line"><a name="l00100"></a><span class="lineno"> 100</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00101"></a><span class="lineno"> 101</span>&#160; sums = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>.ptr);</div><div class="line"><a name="l00102"></a><span class="lineno"> 102</span>&#160;</div><div class="line"><a name="l00103"></a><span class="lineno"> 103</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00104"></a><span class="lineno"> 104</span>&#160; normalize_value = (<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))rsqrt(fmax(sums, <a class="code" href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a>));</div><div class="line"><a name="l00105"></a><span class="lineno"> 105</span>&#160;</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span>&#160; vstore16(in * normalize_value, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>.ptr);</div><div class="line"><a name="l00107"></a><span class="lineno"> 107</span>&#160;}</div><div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_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#l00113">helpers.h:113</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000410<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>
Jenkins514be652019-02-28 12:25:18 +0000411<div class="ttc" id="reduction__operation_8cl_xhtml_ab0df00f5333da51860deb93deb44a782"><div class="ttname"><a href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a></div><div class="ttdeci">DATA_TYPE sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate sum of a vector.</div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00052">reduction_operation.cl:52</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000412<div class="ttc" id="_asymm_helpers_8cpp_xhtml_a552dc3787d7ea1675f3e4e8993501d58"><div class="ttname"><a href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a></div><div class="ttdeci">constexpr float epsilon</div><div class="ttdef"><b>Definition:</b> <a href="_asymm_helpers_8cpp_source.xhtml#l00033">AsymmHelpers.cpp:33</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000413<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>
414<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ae9d2dc29c2789c253406f9b304cc75a8"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
415<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#l00151">helpers.h:151</a></div></div>
416<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#l00057">helpers.h:57</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000417</div><!-- fragment -->
Jenkins514be652019-02-28 12:25:18 +0000418<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00113">CONVERT_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="_asymm_helpers_8cpp_source.xhtml#l00033">epsilon</a>, <a class="el" href="_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00057">VEC_DATA_TYPE</a>.</p>
419
Jenkinsb9abeae2018-11-22 11:58:08 +0000420</div>
421</div>
422<a id="a809db37fce4c2e573b09b0af0d3b95ee"></a>
423<h2 class="memtitle"><span class="permalink"><a href="#a809db37fce4c2e573b09b0af0d3b95ee">&#9670;&nbsp;</a></span>l2_normalize_z()</h2>
424
425<div class="memitem">
426<div class="memproto">
427 <table class="memname">
428 <tr>
429 <td class="memname">__kernel void l2_normalize_z </td>
430 <td>(</td>
431 <td class="paramtype">__global uchar *&#160;</td>
432 <td class="paramname"><em>src_ptr</em>, </td>
433 </tr>
434 <tr>
435 <td class="paramkey"></td>
436 <td></td>
437 <td class="paramtype">uint&#160;</td>
438 <td class="paramname"><em>src_stride_x</em>, </td>
439 </tr>
440 <tr>
441 <td class="paramkey"></td>
442 <td></td>
443 <td class="paramtype">uint&#160;</td>
444 <td class="paramname"><em>src_step_x</em>, </td>
445 </tr>
446 <tr>
447 <td class="paramkey"></td>
448 <td></td>
449 <td class="paramtype">uint&#160;</td>
450 <td class="paramname"><em>src_stride_y</em>, </td>
451 </tr>
452 <tr>
453 <td class="paramkey"></td>
454 <td></td>
455 <td class="paramtype">uint&#160;</td>
456 <td class="paramname"><em>src_step_y</em>, </td>
457 </tr>
458 <tr>
459 <td class="paramkey"></td>
460 <td></td>
461 <td class="paramtype">uint&#160;</td>
462 <td class="paramname"><em>src_stride_z</em>, </td>
463 </tr>
464 <tr>
465 <td class="paramkey"></td>
466 <td></td>
467 <td class="paramtype">uint&#160;</td>
468 <td class="paramname"><em>src_step_z</em>, </td>
469 </tr>
470 <tr>
471 <td class="paramkey"></td>
472 <td></td>
473 <td class="paramtype">uint&#160;</td>
474 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
475 </tr>
476 <tr>
477 <td class="paramkey"></td>
478 <td></td>
479 <td class="paramtype">__global uchar *&#160;</td>
480 <td class="paramname"><em>sum_ptr</em>, </td>
481 </tr>
482 <tr>
483 <td class="paramkey"></td>
484 <td></td>
485 <td class="paramtype">uint&#160;</td>
486 <td class="paramname"><em>sum_stride_x</em>, </td>
487 </tr>
488 <tr>
489 <td class="paramkey"></td>
490 <td></td>
491 <td class="paramtype">uint&#160;</td>
492 <td class="paramname"><em>sum_step_x</em>, </td>
493 </tr>
494 <tr>
495 <td class="paramkey"></td>
496 <td></td>
497 <td class="paramtype">uint&#160;</td>
498 <td class="paramname"><em>sum_stride_y</em>, </td>
499 </tr>
500 <tr>
501 <td class="paramkey"></td>
502 <td></td>
503 <td class="paramtype">uint&#160;</td>
504 <td class="paramname"><em>sum_step_y</em>, </td>
505 </tr>
506 <tr>
507 <td class="paramkey"></td>
508 <td></td>
509 <td class="paramtype">uint&#160;</td>
510 <td class="paramname"><em>sum_stride_z</em>, </td>
511 </tr>
512 <tr>
513 <td class="paramkey"></td>
514 <td></td>
515 <td class="paramtype">uint&#160;</td>
516 <td class="paramname"><em>sum_step_z</em>, </td>
517 </tr>
518 <tr>
519 <td class="paramkey"></td>
520 <td></td>
521 <td class="paramtype">uint&#160;</td>
522 <td class="paramname"><em>sum_offset_first_element_in_bytes</em>, </td>
523 </tr>
524 <tr>
525 <td class="paramkey"></td>
526 <td></td>
527 <td class="paramtype">__global uchar *&#160;</td>
528 <td class="paramname"><em>dst_ptr</em>, </td>
529 </tr>
530 <tr>
531 <td class="paramkey"></td>
532 <td></td>
533 <td class="paramtype">uint&#160;</td>
534 <td class="paramname"><em>dst_stride_x</em>, </td>
535 </tr>
536 <tr>
537 <td class="paramkey"></td>
538 <td></td>
539 <td class="paramtype">uint&#160;</td>
540 <td class="paramname"><em>dst_step_x</em>, </td>
541 </tr>
542 <tr>
543 <td class="paramkey"></td>
544 <td></td>
545 <td class="paramtype">uint&#160;</td>
546 <td class="paramname"><em>dst_stride_y</em>, </td>
547 </tr>
548 <tr>
549 <td class="paramkey"></td>
550 <td></td>
551 <td class="paramtype">uint&#160;</td>
552 <td class="paramname"><em>dst_step_y</em>, </td>
553 </tr>
554 <tr>
555 <td class="paramkey"></td>
556 <td></td>
557 <td class="paramtype">uint&#160;</td>
558 <td class="paramname"><em>dst_stride_z</em>, </td>
559 </tr>
560 <tr>
561 <td class="paramkey"></td>
562 <td></td>
563 <td class="paramtype">uint&#160;</td>
564 <td class="paramname"><em>dst_step_z</em>, </td>
565 </tr>
566 <tr>
567 <td class="paramkey"></td>
568 <td></td>
569 <td class="paramtype">uint&#160;</td>
570 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
571 </tr>
572 <tr>
573 <td class="paramkey"></td>
574 <td></td>
575 <td class="paramtype"><a class="el" href="convolution9x9_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>&#160;</td>
576 <td class="paramname"><em>epsilon</em>&#160;</td>
577 </tr>
578 <tr>
579 <td></td>
580 <td>)</td>
581 <td></td><td></td>
582 </tr>
583 </table>
584</div><div class="memdoc">
585
586<p>This kernel performs l2 normalization on z-axis. </p>
587<dl class="section note"><dt>Note</dt><dd>The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float </dd>
588<dd>
589The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32</dd></dl>
590<dl class="params"><dt>Parameters</dt><dd>
591 <table class="params">
592 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: F16/F32 </td></tr>
593 <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>
594 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_x</td><td>src_stride_x * number of elements along Y processed per workitem(in bytes) </td></tr>
595 <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>
596 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_y</td><td>src_stride_y * number of elements along X processed per workitem(in bytes) </td></tr>
597 <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>
598 <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>
599 <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>
600 <tr><td class="paramdir">[in]</td><td class="paramname">sum_ptr</td><td>Pointer to the source tensor. Supported data types: F16/F32 </td></tr>
601 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_x</td><td>Stride of the source tensor in X dimension (in bytes) </td></tr>
602 <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>
603 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_y</td><td>Stride of the source tensor in Y dimension (in bytes) </td></tr>
604 <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>
605 <tr><td class="paramdir">[in]</td><td class="paramname">sum_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
606 <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>
607 <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 source tensor </td></tr>
608 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: same as <code>src_ptr</code> </td></tr>
609 <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>
610 <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>
611 <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>
612 <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>
613 <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>
614 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
615 <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>
616 <tr><td class="paramdir">[in]</td><td class="paramname">epsilon</td><td>Epsilon value </td></tr>
617 </table>
618 </dd>
619</dl>
620
621<p class="definition">Definition at line <a class="el" href="l2__normalize_8cl_source.xhtml#l00139">139</a> of file <a class="el" href="l2__normalize_8cl_source.xhtml">l2_normalize.cl</a>.</p>
Jenkins514be652019-02-28 12:25:18 +0000622<div class="fragment"><div class="line"><a name="l00144"></a><span class="lineno"> 144</span>&#160;{</div><div class="line"><a name="l00145"></a><span class="lineno"> 145</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a> = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a>);</div><div class="line"><a name="l00146"></a><span class="lineno"> 146</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</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#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>);</div><div class="line"><a name="l00147"></a><span class="lineno"> 147</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</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#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>);</div><div class="line"><a name="l00148"></a><span class="lineno"> 148</span>&#160;</div><div class="line"><a name="l00149"></a><span class="lineno"> 149</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00150"></a><span class="lineno"> 150</span>&#160; in = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">src</a>.ptr);</div><div class="line"><a name="l00151"></a><span class="lineno"> 151</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span>&#160; sums = vload16(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a>.ptr);</div><div class="line"><a name="l00153"></a><span class="lineno"> 153</span>&#160;</div><div class="line"><a name="l00154"></a><span class="lineno"> 154</span>&#160; <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, 16)</div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span>&#160; normalize_value = (<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))rsqrt(fmax(sums, <a class="code" href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a>));</div><div class="line"><a name="l00156"></a><span class="lineno"> 156</span>&#160;</div><div class="line"><a name="l00157"></a><span class="lineno"> 157</span>&#160; vstore16(in * normalize_value, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a>.ptr);</div><div class="line"><a name="l00158"></a><span class="lineno"> 158</span>&#160;}</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>
623<div class="ttc" id="reduction__operation_8cl_xhtml_ab0df00f5333da51860deb93deb44a782"><div class="ttname"><a href="reduction__operation_8cl.xhtml#ab0df00f5333da51860deb93deb44a782">sum</a></div><div class="ttdeci">DATA_TYPE sum(__global const DATA_TYPE *input)</div><div class="ttdoc">Calculate sum of a vector.</div><div class="ttdef"><b>Definition:</b> <a href="reduction__operation_8cl_source.xhtml#l00052">reduction_operation.cl:52</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000624<div class="ttc" id="_asymm_helpers_8cpp_xhtml_a552dc3787d7ea1675f3e4e8993501d58"><div class="ttname"><a href="_asymm_helpers_8cpp.xhtml#a552dc3787d7ea1675f3e4e8993501d58">epsilon</a></div><div class="ttdeci">constexpr float epsilon</div><div class="ttdef"><b>Definition:</b> <a href="_asymm_helpers_8cpp_source.xhtml#l00033">AsymmHelpers.cpp:33</a></div></div>
Jenkins514be652019-02-28 12:25:18 +0000625<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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00160">helpers.h:160</a></div></div>
626<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>
627<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_ae9d2dc29c2789c253406f9b304cc75a8"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#ae9d2dc29c2789c253406f9b304cc75a8">arm_compute::test::validation::src</a></div><div class="ttdeci">cast configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_cast_8cpp_source.xhtml#l00169">Cast.cpp:169</a></div></div>
628<div class="ttc" id="src_2core_2_c_l_2cl__kernels_2_helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="src_2core_2_c_l_2cl__kernels_2_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="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00128">helpers.h:128</a></div></div>
629<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#l00057">helpers.h:57</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100630</div><!-- fragment -->
Jenkins514be652019-02-28 12:25:18 +0000631<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00128">CONVERT_TO_TENSOR3D_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="_asymm_helpers_8cpp_source.xhtml#l00033">epsilon</a>, <a class="el" href="_cast_8cpp_source.xhtml#l00169">arm_compute::test::validation::src</a>, <a class="el" href="reduction__operation_8cl_source.xhtml#l00052">sum()</a>, and <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00057">VEC_DATA_TYPE</a>.</p>
632
Kaizen8938bd32017-09-28 14:38:23 +0100633</div>
634</div>
635</div><!-- contents -->
636</div><!-- doc-content -->
637<!-- start footer part -->
638<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
639 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000640 <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="l2__normalize_8cl.xhtml">l2_normalize.cl</a></li>
Jenkins514be652019-02-28 12:25:18 +0000641 <li class="footer">Generated on Thu Feb 28 2019 12:24:56 for Compute Library by
Kaizen8938bd32017-09-28 14:38:23 +0100642 <a href="http://www.doxygen.org/index.html">
Jenkins514be652019-02-28 12:25:18 +0000643 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.15 </li>
Kaizen8938bd32017-09-28 14:38:23 +0100644 </ul>
645</div>
646</body>
647</html>