blob: 8ccd7bcf529e99ada38a9b12471081aa63d4828a [file] [log] [blame]
Kaizen8938bd32017-09-28 14:38:23 +01001<!-- 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"/>
Jenkinsb9abeae2018-11-22 11:58:08 +00007<meta name="generator" content="Doxygen 1.8.13"/>
Kaizen8938bd32017-09-28 14:38:23 +01008<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
9<title>Compute Library: src/core/CL/cl_kernels/quantization_layer.cl File Reference</title>
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 +000015<script type="text/javascript" src="navtreedata.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010016<script type="text/javascript" src="navtree.js"></script>
17<script type="text/javascript">
18 $(document).ready(initResizable);
Kaizen8938bd32017-09-28 14:38:23 +010019</script>
20<link href="search/search.css" rel="stylesheet" type="text/css"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +000021<script type="text/javascript" src="search/searchdata.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010022<script type="text/javascript" src="search/search.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010023<script type="text/x-mathjax-config">
24 MathJax.Hub.Config({
25 extensions: ["tex2jax.js"],
26 jax: ["input/TeX","output/HTML-CSS"],
27});
Anthony Barbier8140e1e2017-12-14 23:48:46 +000028</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +010029<link href="doxygen.css" rel="stylesheet" type="text/css" />
30</head>
31<body>
32<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
33<div id="titlearea">
34<table cellspacing="0" cellpadding="0">
35 <tbody>
36 <tr style="height: 56px;">
37 <td style="padding-left: 0.5em;">
38 <div id="projectname">Compute Library
Jenkinsb9abeae2018-11-22 11:58:08 +000039 &#160;<span id="projectnumber">18.11</span>
Kaizen8938bd32017-09-28 14:38:23 +010040 </div>
41 </td>
42 </tr>
43 </tbody>
44</table>
45</div>
46<!-- end header part -->
Jenkinsb9abeae2018-11-22 11:58:08 +000047<!-- Generated by Doxygen 1.8.13 -->
Kaizen8938bd32017-09-28 14:38:23 +010048<script type="text/javascript">
49var searchBox = new SearchBox("searchBox", "search",false,'Search');
50</script>
Jenkinsb9abeae2018-11-22 11:58:08 +000051<script type="text/javascript" src="menudata.js"></script>
52<script type="text/javascript" src="menu.js"></script>
53<script type="text/javascript">
54$(function() {
55 initMenu('',true,false,'search.php','Search');
56 $(document).ready(function() { init_search(); });
57});
58</script>
59<div id="main-nav"></div>
Kaizen8938bd32017-09-28 14:38:23 +010060</div><!-- top -->
61<div id="side-nav" class="ui-resizable side-nav-resizable">
62 <div id="nav-tree">
63 <div id="nav-tree-contents">
64 <div id="nav-sync" class="sync"></div>
65 </div>
66 </div>
67 <div id="splitbar" style="-moz-user-select:none;"
68 class="ui-resizable-handle">
69 </div>
70</div>
71<script type="text/javascript">
72$(document).ready(function(){initNavTree('quantization__layer_8cl.xhtml','');});
73</script>
74<div id="doc-content">
75<!-- window showing the filter options -->
76<div id="MSearchSelectWindow"
77 onmouseover="return searchBox.OnSearchSelectShow()"
78 onmouseout="return searchBox.OnSearchSelectHide()"
79 onkeydown="return searchBox.OnSearchSelectKey(event)">
Anthony Barbier8140e1e2017-12-14 23:48:46 +000080</div>
Kaizen8938bd32017-09-28 14:38:23 +010081
82<!-- iframe showing the search results (closed by default) -->
83<div id="MSearchResultsWindow">
84<iframe src="javascript:void(0)" frameborder="0"
85 name="MSearchResults" id="MSearchResults">
86</iframe>
87</div>
88
89<div class="header">
90 <div class="summary">
91<a href="#func-members">Functions</a> </div>
92 <div class="headertitle">
93<div class="title">quantization_layer.cl File Reference</div> </div>
94</div><!--header-->
95<div class="contents">
Jenkinsb9abeae2018-11-22 11:58:08 +000096<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 +010097</div>
98<p><a href="quantization__layer_8cl_source.xhtml">Go to the source code of this file.</a></p>
99<table class="memberdecls">
100<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
101Functions</h2></td></tr>
102<tr class="memitem:a1f16fa20ae72c76df16155e6a3ad20c5"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="quantization__layer_8cl.xhtml#a1f16fa20ae72c76df16155e6a3ad20c5">quantization_layer</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, __global uchar *min_max_ptr, uint min_max_stride_x, uint min_max_step_x, uint min_max_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000103<tr class="memdesc:a1f16fa20ae72c76df16155e6a3ad20c5"><td class="mdescLeft">&#160;</td><td class="mdescRight">This performs the quantization of floating point inputs to 8-bit unsigned integers. <a href="#a1f16fa20ae72c76df16155e6a3ad20c5">More...</a><br /></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100104<tr class="separator:a1f16fa20ae72c76df16155e6a3ad20c5"><td class="memSeparator" colspan="2">&#160;</td></tr>
105</table>
106<h2 class="groupheader">Function Documentation</h2>
Jenkinsb9abeae2018-11-22 11:58:08 +0000107<a id="a1f16fa20ae72c76df16155e6a3ad20c5"></a>
108<h2 class="memtitle"><span class="permalink"><a href="#a1f16fa20ae72c76df16155e6a3ad20c5">&#9670;&nbsp;</a></span>quantization_layer()</h2>
109
Kaizen8938bd32017-09-28 14:38:23 +0100110<div class="memitem">
111<div class="memproto">
112 <table class="memname">
113 <tr>
114 <td class="memname">__kernel void quantization_layer </td>
115 <td>(</td>
116 <td class="paramtype">__global uchar *&#160;</td>
117 <td class="paramname"><em>input_ptr</em>, </td>
118 </tr>
119 <tr>
120 <td class="paramkey"></td>
121 <td></td>
122 <td class="paramtype">uint&#160;</td>
123 <td class="paramname"><em>input_stride_x</em>, </td>
124 </tr>
125 <tr>
126 <td class="paramkey"></td>
127 <td></td>
128 <td class="paramtype">uint&#160;</td>
129 <td class="paramname"><em>input_step_x</em>, </td>
130 </tr>
131 <tr>
132 <td class="paramkey"></td>
133 <td></td>
134 <td class="paramtype">uint&#160;</td>
135 <td class="paramname"><em>input_stride_y</em>, </td>
136 </tr>
137 <tr>
138 <td class="paramkey"></td>
139 <td></td>
140 <td class="paramtype">uint&#160;</td>
141 <td class="paramname"><em>input_step_y</em>, </td>
142 </tr>
143 <tr>
144 <td class="paramkey"></td>
145 <td></td>
146 <td class="paramtype">uint&#160;</td>
147 <td class="paramname"><em>input_stride_z</em>, </td>
148 </tr>
149 <tr>
150 <td class="paramkey"></td>
151 <td></td>
152 <td class="paramtype">uint&#160;</td>
153 <td class="paramname"><em>input_step_z</em>, </td>
154 </tr>
155 <tr>
156 <td class="paramkey"></td>
157 <td></td>
158 <td class="paramtype">uint&#160;</td>
159 <td class="paramname"><em>input_offset_first_element_in_bytes</em>, </td>
160 </tr>
161 <tr>
162 <td class="paramkey"></td>
163 <td></td>
164 <td class="paramtype">__global uchar *&#160;</td>
165 <td class="paramname"><em>output_ptr</em>, </td>
166 </tr>
167 <tr>
168 <td class="paramkey"></td>
169 <td></td>
170 <td class="paramtype">uint&#160;</td>
171 <td class="paramname"><em>output_stride_x</em>, </td>
172 </tr>
173 <tr>
174 <td class="paramkey"></td>
175 <td></td>
176 <td class="paramtype">uint&#160;</td>
177 <td class="paramname"><em>output_step_x</em>, </td>
178 </tr>
179 <tr>
180 <td class="paramkey"></td>
181 <td></td>
182 <td class="paramtype">uint&#160;</td>
183 <td class="paramname"><em>output_stride_y</em>, </td>
184 </tr>
185 <tr>
186 <td class="paramkey"></td>
187 <td></td>
188 <td class="paramtype">uint&#160;</td>
189 <td class="paramname"><em>output_step_y</em>, </td>
190 </tr>
191 <tr>
192 <td class="paramkey"></td>
193 <td></td>
194 <td class="paramtype">uint&#160;</td>
195 <td class="paramname"><em>output_stride_z</em>, </td>
196 </tr>
197 <tr>
198 <td class="paramkey"></td>
199 <td></td>
200 <td class="paramtype">uint&#160;</td>
201 <td class="paramname"><em>output_step_z</em>, </td>
202 </tr>
203 <tr>
204 <td class="paramkey"></td>
205 <td></td>
206 <td class="paramtype">uint&#160;</td>
207 <td class="paramname"><em>output_offset_first_element_in_bytes</em>, </td>
208 </tr>
209 <tr>
210 <td class="paramkey"></td>
211 <td></td>
212 <td class="paramtype">__global uchar *&#160;</td>
213 <td class="paramname"><em>min_max_ptr</em>, </td>
214 </tr>
215 <tr>
216 <td class="paramkey"></td>
217 <td></td>
218 <td class="paramtype">uint&#160;</td>
219 <td class="paramname"><em>min_max_stride_x</em>, </td>
220 </tr>
221 <tr>
222 <td class="paramkey"></td>
223 <td></td>
224 <td class="paramtype">uint&#160;</td>
225 <td class="paramname"><em>min_max_step_x</em>, </td>
226 </tr>
227 <tr>
228 <td class="paramkey"></td>
229 <td></td>
230 <td class="paramtype">uint&#160;</td>
231 <td class="paramname"><em>min_max_offset_first_element_in_bytes</em>&#160;</td>
232 </tr>
233 <tr>
234 <td></td>
235 <td>)</td>
236 <td></td><td></td>
237 </tr>
238 </table>
239</div><div class="memdoc">
240
241<p>This performs the quantization of floating point inputs to 8-bit unsigned integers. </p>
242<dl class="params"><dt>Parameters</dt><dd>
243 <table class="params">
244 <tr><td class="paramdir">[in]</td><td class="paramname">input_ptr</td><td>Pointer to the source image. Supported data types: F32 </td></tr>
245 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
246 <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>
247 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
248 <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>
249 <tr><td class="paramdir">[in]</td><td class="paramname">input_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
250 <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>
251 <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 source image </td></tr>
252 <tr><td class="paramdir">[out]</td><td class="paramname">output_ptr</td><td>Pointer to the destination image. Supported data types: U8 </td></tr>
253 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_x</td><td>Stride of the destination image in X dimension (in bytes) </td></tr>
254 <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>
255 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_y</td><td>Stride of the destination image in Y dimension (in bytes) </td></tr>
256 <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>
257 <tr><td class="paramdir">[in]</td><td class="paramname">output_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
258 <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>
259 <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 image </td></tr>
260 <tr><td class="paramdir">[in]</td><td class="paramname">min_max_ptr</td><td>Pointer to the min/max vector. Minimum value in position 0, maximum value in position 1. Supported data types: F32. </td></tr>
261 <tr><td class="paramdir">[in]</td><td class="paramname">min_max_stride_x</td><td>Stride of the min/max vector in X dimension (in bytes) </td></tr>
262 <tr><td class="paramdir">[in]</td><td class="paramname">min_max_step_x</td><td>min_max_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
263 <tr><td class="paramdir">[in]</td><td class="paramname">min_max_offset_first_element_in_bytes</td><td>The offset of the first element in the min/max vector </td></tr>
264 </table>
265 </dd>
266</dl>
267
Jenkinsb9abeae2018-11-22 11:58:08 +0000268<p class="definition">Definition at line <a class="el" href="quantization__layer_8cl_source.xhtml#l00049">49</a> of file <a class="el" href="quantization__layer_8cl_source.xhtml">quantization_layer.cl</a>.</p>
Kaizen8938bd32017-09-28 14:38:23 +0100269
Jenkinsb9abeae2018-11-22 11:58:08 +0000270<p class="reference">References <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00125">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00159">Tensor3D::ptr</a>, and <a class="el" href="min_8h_source.xhtml#l00039">arm_compute::wrapper::vmin()</a>.</p>
271<div class="fragment"><div class="line"><a name="l00053"></a><span class="lineno"> 53</span>&#160;{</div><div class="line"><a name="l00054"></a><span class="lineno"> 54</span>&#160; <span class="comment">// Get pixels pointer</span></div><div class="line"><a name="l00055"></a><span class="lineno"> 55</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> input = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(input);</div><div class="line"><a name="l00056"></a><span class="lineno"> 56</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> output = <a class="code" href="src_2core_2_c_l_2cl__kernels_2_helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(output);</div><div class="line"><a name="l00057"></a><span class="lineno"> 57</span>&#160;</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>&#160; <span class="comment">// min_max_value.s0 = min, min_max_value.s1 = max</span></div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>&#160; <span class="keyword">const</span> float2 min_max_value = vload2(0, (__global <span class="keywordtype">float</span> *)(min_max_ptr + min_max_offset_first_element_in_bytes));</div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span>&#160;</div><div class="line"><a name="l00061"></a><span class="lineno"> 61</span>&#160; <span class="keyword">const</span> float4 <a class="code" href="namespacearm__compute_1_1wrapper.xhtml#ab026a1aeec48074092f548c130c323f4">vmin</a> = (float4)min_max_value.s0;</div><div class="line"><a name="l00062"></a><span class="lineno"> 62</span>&#160; <span class="keyword">const</span> float4 vrange = (float4)(min_max_value.s1 - min_max_value.s0);</div><div class="line"><a name="l00063"></a><span class="lineno"> 63</span>&#160;</div><div class="line"><a name="l00064"></a><span class="lineno"> 64</span>&#160; <span class="comment">// Load data</span></div><div class="line"><a name="l00065"></a><span class="lineno"> 65</span>&#160; float4 data = vload4(0, (__global <span class="keywordtype">float</span> *)input.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00066"></a><span class="lineno"> 66</span>&#160;</div><div class="line"><a name="l00067"></a><span class="lineno"> 67</span>&#160; <span class="comment">// Map float values to range [0.0, 1.0]</span></div><div class="line"><a name="l00068"></a><span class="lineno"> 68</span>&#160; data = (data - <a class="code" href="namespacearm__compute_1_1wrapper.xhtml#ab026a1aeec48074092f548c130c323f4">vmin</a>) / vrange;</div><div class="line"><a name="l00069"></a><span class="lineno"> 69</span>&#160;</div><div class="line"><a name="l00070"></a><span class="lineno"> 70</span>&#160; <span class="comment">// Quantize and saturate</span></div><div class="line"><a name="l00071"></a><span class="lineno"> 71</span>&#160; uchar4 res = convert_uchar4_sat(data * 256.0f);</div><div class="line"><a name="l00072"></a><span class="lineno"> 72</span>&#160;</div><div class="line"><a name="l00073"></a><span class="lineno"> 73</span>&#160; <span class="comment">// Store result</span></div><div class="line"><a name="l00074"></a><span class="lineno"> 74</span>&#160; vstore4(res, 0, (__global uchar *)output.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00075"></a><span class="lineno"> 75</span>&#160;}</div><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#l00157">helpers.h:157</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000272<div class="ttc" id="namespacearm__compute_1_1wrapper_xhtml_ab026a1aeec48074092f548c130c323f4"><div class="ttname"><a href="namespacearm__compute_1_1wrapper.xhtml#ab026a1aeec48074092f548c130c323f4">arm_compute::wrapper::vmin</a></div><div class="ttdeci">uint8x8_t vmin(const uint8x8_t &amp;a, const uint8x8_t &amp;b)</div><div class="ttdef"><b>Definition:</b> <a href="min_8h_source.xhtml#l00039">min.h:39</a></div></div>
Jenkinsb9abeae2018-11-22 11:58:08 +0000273<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#l00125">helpers.h:125</a></div></div>
274<div class="ttc" id="struct_tensor3_d_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">Tensor3D::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="src_2core_2_c_l_2cl__kernels_2_helpers_8h_source.xhtml#l00159">helpers.h:159</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100275</div><!-- fragment -->
276</div>
277</div>
278</div><!-- contents -->
279</div><!-- doc-content -->
280<!-- start footer part -->
281<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
282 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000283 <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="quantization__layer_8cl.xhtml">quantization_layer.cl</a></li>
Jenkinsb9abeae2018-11-22 11:58:08 +0000284 <li class="footer">Generated on Thu Nov 22 2018 11:57:43 for Compute Library by
Kaizen8938bd32017-09-28 14:38:23 +0100285 <a href="http://www.doxygen.org/index.html">
Jenkinsb9abeae2018-11-22 11:58:08 +0000286 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.13 </li>
Kaizen8938bd32017-09-28 14:38:23 +0100287 </ul>
288</div>
289</body>
290</html>