Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 06ea048f062a50404b1b3998a61a45449c2d1f0f / . / documentation / depthwise__convolution_8cl.xhtml
blob: 0a7875d8cf8e79369ae5f877c5fabd3a9ce66672 [file] [log] [blame]
Kaizen8938bd32017-09-28 14:38:23 +0100[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"/>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]8<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
9<title>Compute Library: src/core/CL/cl_kernels/depthwise_convolution.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 +0000[diff] [blame]15<script type="text/javascript" src="navtreedata.js"></script>
Kaizen8938bd32017-09-28 14:38:23 +0100[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>
Kaizen8938bd32017-09-28 14:38:23 +0100[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(); });
Kaizen8938bd32017-09-28 14:38:23 +0100[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>
Kaizen8938bd32017-09-28 14:38:23 +0100[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;">
42 <div id="projectname">Compute Library
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame^]43 &#160;<span id="projectnumber">18.02</span>
Kaizen8938bd32017-09-28 14:38:23 +0100[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 -->
Kaizen8938bd32017-09-28 14:38:23 +0100[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>
58 <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
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('depthwise__convolution_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>
Kaizen8938bd32017-09-28 14:38:23 +0100[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">
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame^]116 <div class="summary">
117<a href="#define-members">Macros</a> &#124;
118<a href="#func-members">Functions</a> </div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]119 <div class="headertitle">
120<div class="title">depthwise_convolution.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 />
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]124</div>
125<p><a href="depthwise__convolution_8cl_source.xhtml">Go to the source code of this file.</a></p>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame^]126<table class="memberdecls">
127<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
128Macros</h2></td></tr>
129<tr class="memitem:aa5ab4cf85fe56de3821727cfcef327d9"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(acc, src0, weights_row0)</td></tr>
130<tr class="separator:aa5ab4cf85fe56de3821727cfcef327d9"><td class="memSeparator" colspan="2">&#160;</td></tr>
131<tr class="memitem:a31cbdd6b28ee6331839432f69d87fdcb"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(acc, src0, src1, weights_row0)</td></tr>
132<tr class="separator:a31cbdd6b28ee6331839432f69d87fdcb"><td class="memSeparator" colspan="2">&#160;</td></tr>
133</table><table class="memberdecls">
134<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
135Functions</h2></td></tr>
136<tr class="memitem:a2fbbf5b73260b66b310d4d857db5289e"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a2fbbf5b73260b66b310d4d857db5289e">depthwise_convolution_3x3_stridex1_stridey1_bifrost</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 *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes, __global uchar *weights_ptr, uint weights_stride_x, uint weights_step_x, uint weights_stride_y, uint weights_step_y, uint weights_stride_z, uint weights_step_z, uint weights_offset_first_element_in_bytes, __global uchar *biases_ptr, uint biases_stride_x, uint biases_step_x, uint biases_offset_first_element_in_bytes)</td></tr>
137<tr class="memdesc:a2fbbf5b73260b66b310d4d857db5289e"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 1. <a href="#a2fbbf5b73260b66b310d4d857db5289e">More...</a><br /></td></tr>
138<tr class="separator:a2fbbf5b73260b66b310d4d857db5289e"><td class="memSeparator" colspan="2">&#160;</td></tr>
139<tr class="memitem:a328ab548c536089e8238fb33de617bec"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a328ab548c536089e8238fb33de617bec">depthwise_convolution_3x3_stridex2_stridey2_bifrost</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 *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes, __global uchar *weights_ptr, uint weights_stride_x, uint weights_step_x, uint weights_stride_y, uint weights_step_y, uint weights_stride_z, uint weights_step_z, uint weights_offset_first_element_in_bytes, __global uchar *biases_ptr, uint biases_stride_x, uint biases_step_x, uint biases_offset_first_element_in_bytes)</td></tr>
140<tr class="memdesc:a328ab548c536089e8238fb33de617bec"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 2. <a href="#a328ab548c536089e8238fb33de617bec">More...</a><br /></td></tr>
141<tr class="separator:a328ab548c536089e8238fb33de617bec"><td class="memSeparator" colspan="2">&#160;</td></tr>
142</table>
143<h2 class="groupheader">Macro Definition Documentation</h2>
144<a class="anchor" id="aa5ab4cf85fe56de3821727cfcef327d9"></a>
145<div class="memitem">
146<div class="memproto">
147 <table class="memname">
148 <tr>
149 <td class="memname">#define CONVOLUTION1x3_BIFROST2X1_STRIDE1</td>
150 <td>(</td>
151 <td class="paramtype">&#160;</td>
152 <td class="paramname">acc, </td>
153 </tr>
154 <tr>
155 <td class="paramkey"></td>
156 <td></td>
157 <td class="paramtype">&#160;</td>
158 <td class="paramname">src0, </td>
159 </tr>
160 <tr>
161 <td class="paramkey"></td>
162 <td></td>
163 <td class="paramtype">&#160;</td>
164 <td class="paramname">weights_row0&#160;</td>
165 </tr>
166 <tr>
167 <td></td>
168 <td>)</td>
169 <td></td><td></td>
170 </tr>
171 </table>
172</div><div class="memdoc">
173<b>Value:</b><div class="fragment"><div class="line">({ \</div><div class="line"> acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0); \</div><div class="line"> acc.s1 = fma(src0.s1, weights_row0.s0, acc.s1); \</div><div class="line"> acc.s1 = fma(src0.s2, weights_row0.s1, acc.s1); \</div><div class="line"> acc.s1 = fma(src0.s3, weights_row0.s2, acc.s1); \</div><div class="line"> })</div></div><!-- fragment -->
174<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00211">211</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
175
176<p>Referenced by <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00263">depthwise_convolution_3x3_stridex1_stridey1_bifrost()</a>.</p>
177
178</div>
179</div>
180<a class="anchor" id="a31cbdd6b28ee6331839432f69d87fdcb"></a>
181<div class="memitem">
182<div class="memproto">
183 <table class="memname">
184 <tr>
185 <td class="memname">#define CONVOLUTION1x3_BIFROST2X1_STRIDE2</td>
186 <td>(</td>
187 <td class="paramtype">&#160;</td>
188 <td class="paramname">acc, </td>
189 </tr>
190 <tr>
191 <td class="paramkey"></td>
192 <td></td>
193 <td class="paramtype">&#160;</td>
194 <td class="paramname">src0, </td>
195 </tr>
196 <tr>
197 <td class="paramkey"></td>
198 <td></td>
199 <td class="paramtype">&#160;</td>
200 <td class="paramname">src1, </td>
201 </tr>
202 <tr>
203 <td class="paramkey"></td>
204 <td></td>
205 <td class="paramtype">&#160;</td>
206 <td class="paramname">weights_row0&#160;</td>
207 </tr>
208 <tr>
209 <td></td>
210 <td>)</td>
211 <td></td><td></td>
212 </tr>
213 </table>
214</div><div class="memdoc">
215<b>Value:</b><div class="fragment"><div class="line">({ \</div><div class="line"> acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0); \</div><div class="line"> acc.s1 = fma(src0.s2, weights_row0.s0, acc.s1); \</div><div class="line"> acc.s1 = fma(src0.s3, weights_row0.s1, acc.s1); \</div><div class="line"> acc.s1 = fma(src1.s0, weights_row0.s2, acc.s1); \</div><div class="line"> })</div></div><!-- fragment -->
216<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00221">221</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
217
218<p>Referenced by <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00360">depthwise_convolution_3x3_stridex2_stridey2_bifrost()</a>.</p>
219
220</div>
221</div>
222<h2 class="groupheader">Function Documentation</h2>
223<a class="anchor" id="a2fbbf5b73260b66b310d4d857db5289e"></a>
224<div class="memitem">
225<div class="memproto">
226 <table class="memname">
227 <tr>
228 <td class="memname">__kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost </td>
229 <td>(</td>
230 <td class="paramtype">__global uchar *&#160;</td>
231 <td class="paramname"><em>src_ptr</em>, </td>
232 </tr>
233 <tr>
234 <td class="paramkey"></td>
235 <td></td>
236 <td class="paramtype">uint&#160;</td>
237 <td class="paramname"><em>src_stride_x</em>, </td>
238 </tr>
239 <tr>
240 <td class="paramkey"></td>
241 <td></td>
242 <td class="paramtype">uint&#160;</td>
243 <td class="paramname"><em>src_step_x</em>, </td>
244 </tr>
245 <tr>
246 <td class="paramkey"></td>
247 <td></td>
248 <td class="paramtype">uint&#160;</td>
249 <td class="paramname"><em>src_stride_y</em>, </td>
250 </tr>
251 <tr>
252 <td class="paramkey"></td>
253 <td></td>
254 <td class="paramtype">uint&#160;</td>
255 <td class="paramname"><em>src_step_y</em>, </td>
256 </tr>
257 <tr>
258 <td class="paramkey"></td>
259 <td></td>
260 <td class="paramtype">uint&#160;</td>
261 <td class="paramname"><em>src_stride_z</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_step_z</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_offset_first_element_in_bytes</em>, </td>
274 </tr>
275 <tr>
276 <td class="paramkey"></td>
277 <td></td>
278 <td class="paramtype">__global uchar *&#160;</td>
279 <td class="paramname"><em>dst_ptr</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>dst_stride_x</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>dst_step_x</em>, </td>
292 </tr>
293 <tr>
294 <td class="paramkey"></td>
295 <td></td>
296 <td class="paramtype">uint&#160;</td>
297 <td class="paramname"><em>dst_stride_y</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>dst_step_y</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>dst_stride_z</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>dst_step_z</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>dst_offset_first_element_in_bytes</em>, </td>
322 </tr>
323 <tr>
324 <td class="paramkey"></td>
325 <td></td>
326 <td class="paramtype">__global uchar *&#160;</td>
327 <td class="paramname"><em>weights_ptr</em>, </td>
328 </tr>
329 <tr>
330 <td class="paramkey"></td>
331 <td></td>
332 <td class="paramtype">uint&#160;</td>
333 <td class="paramname"><em>weights_stride_x</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>weights_step_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>weights_stride_y</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>weights_step_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>weights_stride_z</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>weights_step_z</em>, </td>
364 </tr>
365 <tr>
366 <td class="paramkey"></td>
367 <td></td>
368 <td class="paramtype">uint&#160;</td>
369 <td class="paramname"><em>weights_offset_first_element_in_bytes</em>, </td>
370 </tr>
371 <tr>
372 <td class="paramkey"></td>
373 <td></td>
374 <td class="paramtype">__global uchar *&#160;</td>
375 <td class="paramname"><em>biases_ptr</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>biases_stride_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>biases_step_x</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>biases_offset_first_element_in_bytes</em>&#160;</td>
394 </tr>
395 <tr>
396 <td></td>
397 <td>)</td>
398 <td></td><td></td>
399 </tr>
400 </table>
401</div><div class="memdoc">
402
403<p>This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 1. </p>
404<dl class="params"><dt>Parameters</dt><dd>
405 <table class="params">
406 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source image. Supported data types: F32 </td></tr>
407 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
408 <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>
409 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
410 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_y</td><td>src_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
411 <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 image </td></tr>
412 <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>
413 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
414 <tr><td class="paramdir">[in]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F32 </td></tr>
415 <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>
416 <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>
417 <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>
418 <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>
419 <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>
420 <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>
421 <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>
422 <tr><td class="paramdir">[in]</td><td class="paramname">weights_ptr</td><td>Pointer to the weights tensor. Supported data types: F32 </td></tr>
423 <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_x</td><td>Stride of the weights tensor in X dimension (in bytes) </td></tr>
424 <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_x</td><td>weights_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
425 <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_y</td><td>Stride of the weights tensor in Y dimension (in bytes) </td></tr>
426 <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_y</td><td>weights_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
427 <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_z</td><td>Stride of the weights tensor in Z dimension (in bytes) </td></tr>
428 <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_z</td><td>weights_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
429 <tr><td class="paramdir">[in]</td><td class="paramname">weights_offset_first_element_in_bytes</td><td>The offset of the first element in the biases vector </td></tr>
430 <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>(Optional) Pointer to the biases vector. Supported data types: F32 </td></tr>
431 <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>(Optional) Stride of the biases vector in X dimension (in bytes) </td></tr>
432 <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>(Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
433 <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>(Optional) The offset of the first element in the biases vector </td></tr>
434 </table>
435 </dd>
436</dl>
437
438<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00263">263</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
439
440<p>References <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00117">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00099">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>, <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00211">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00151">Tensor3D::ptr</a>, and <a class="el" href="helpers_8h_source.xhtml#l00290">vector_offset()</a>.</p>
441<div class="fragment"><div class="line"><a name="l00272"></a><span class="lineno"> 272</span>&#160;{</div><div class="line"><a name="l00273"></a><span class="lineno"> 273</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00274"></a><span class="lineno"> 274</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="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00275"></a><span class="lineno"> 275</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> weights = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(weights);</div><div class="line"><a name="l00276"></a><span class="lineno"> 276</span>&#160;</div><div class="line"><a name="l00277"></a><span class="lineno"> 277</span>&#160; float2 pixels0 = 0.0f;</div><div class="line"><a name="l00278"></a><span class="lineno"> 278</span>&#160; float2 pixels1 = 0.0f;</div><div class="line"><a name="l00279"></a><span class="lineno"> 279</span>&#160; float2 pixels2 = 0.0f;</div><div class="line"><a name="l00280"></a><span class="lineno"> 280</span>&#160; float2 pixels3 = 0.0f;</div><div class="line"><a name="l00281"></a><span class="lineno"> 281</span>&#160;</div><div class="line"><a name="l00282"></a><span class="lineno"> 282</span>&#160; __global uchar *weights_addr = (__global uchar *)weights.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>;</div><div class="line"><a name="l00283"></a><span class="lineno"> 283</span>&#160; __global uchar *src_addr = (__global uchar *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0);</div><div class="line"><a name="l00284"></a><span class="lineno"> 284</span>&#160;</div><div class="line"><a name="l00285"></a><span class="lineno"> 285</span>&#160; <span class="comment">// Load the weights</span></div><div class="line"><a name="l00286"></a><span class="lineno"> 286</span>&#160; float3 weights_row0 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 0 * weights_stride_y));</div><div class="line"><a name="l00287"></a><span class="lineno"> 287</span>&#160; float3 weights_row1 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 1 * weights_stride_y));</div><div class="line"><a name="l00288"></a><span class="lineno"> 288</span>&#160; float3 weights_row2 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 2 * weights_stride_y));</div><div class="line"><a name="l00289"></a><span class="lineno"> 289</span>&#160;</div><div class="line"><a name="l00290"></a><span class="lineno"> 290</span>&#160; <span class="comment">// Note: Since each work-item computes 4x2 elements, we need to load 4 rows from the input tensor</span></div><div class="line"><a name="l00291"></a><span class="lineno"> 291</span>&#160; float4 src00 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00292"></a><span class="lineno"> 292</span>&#160; float4 src10 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00293"></a><span class="lineno"> 293</span>&#160; float4 src20 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00294"></a><span class="lineno"> 294</span>&#160; float4 src30 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00295"></a><span class="lineno"> 295</span>&#160; float4 src40 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00296"></a><span class="lineno"> 296</span>&#160; float4 src50 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 5 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00297"></a><span class="lineno"> 297</span>&#160;</div><div class="line"><a name="l00298"></a><span class="lineno"> 298</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src00, weights_row0);</div><div class="line"><a name="l00299"></a><span class="lineno"> 299</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src10, weights_row1);</div><div class="line"><a name="l00300"></a><span class="lineno"> 300</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src20, weights_row2);</div><div class="line"><a name="l00301"></a><span class="lineno"> 301</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src10, weights_row0);</div><div class="line"><a name="l00302"></a><span class="lineno"> 302</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src20, weights_row1);</div><div class="line"><a name="l00303"></a><span class="lineno"> 303</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src30, weights_row2);</div><div class="line"><a name="l00304"></a><span class="lineno"> 304</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src20, weights_row0);</div><div class="line"><a name="l00305"></a><span class="lineno"> 305</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src30, weights_row1);</div><div class="line"><a name="l00306"></a><span class="lineno"> 306</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src40, weights_row2);</div><div class="line"><a name="l00307"></a><span class="lineno"> 307</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src30, weights_row0);</div><div class="line"><a name="l00308"></a><span class="lineno"> 308</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src40, weights_row1);</div><div class="line"><a name="l00309"></a><span class="lineno"> 309</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src50, weights_row2);</div><div class="line"><a name="l00310"></a><span class="lineno"> 310</span>&#160;</div><div class="line"><a name="l00311"></a><span class="lineno"> 311</span>&#160;<span class="preprocessor">#ifdef HAS_BIAS</span></div><div class="line"><a name="l00312"></a><span class="lineno"> 312</span>&#160; <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>(biases);</div><div class="line"><a name="l00313"></a><span class="lineno"> 313</span>&#160;</div><div class="line"><a name="l00314"></a><span class="lineno"> 314</span>&#160; <span class="keywordtype">float</span> bias = *((__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a>(&amp;biases, get_global_id(2))));</div><div class="line"><a name="l00315"></a><span class="lineno"> 315</span>&#160;</div><div class="line"><a name="l00316"></a><span class="lineno"> 316</span>&#160; pixels0 += (float2)bias;</div><div class="line"><a name="l00317"></a><span class="lineno"> 317</span>&#160; pixels1 += (float2)bias;</div><div class="line"><a name="l00318"></a><span class="lineno"> 318</span>&#160; pixels2 += (float2)bias;</div><div class="line"><a name="l00319"></a><span class="lineno"> 319</span>&#160; pixels3 += (float2)bias;</div><div class="line"><a name="l00320"></a><span class="lineno"> 320</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(HAS_BIAS) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00321"></a><span class="lineno"> 321</span>&#160;</div><div class="line"><a name="l00322"></a><span class="lineno"> 322</span>&#160; vstore2(pixels0, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 0 * dst_stride_y));</div><div class="line"><a name="l00323"></a><span class="lineno"> 323</span>&#160; vstore2(pixels1, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 1 * dst_stride_y));</div><div class="line"><a name="l00324"></a><span class="lineno"> 324</span>&#160; vstore2(pixels2, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 2 * dst_stride_y));</div><div class="line"><a name="l00325"></a><span class="lineno"> 325</span>&#160; vstore2(pixels3, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 3 * dst_stride_y));</div><div class="line"><a name="l00326"></a><span class="lineno"> 326</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="helpers_8h_source.xhtml#l00132">helpers.h:132</a></div></div>
442<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#l00117">helpers.h:117</a></div></div>
443<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#l00149">helpers.h:149</a></div></div>
444<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
445<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_convolution_8cpp_source.xhtml#l00137">Convolution.cpp:137</a></div></div>
446<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="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
447<div class="ttc" id="depthwise__convolution_8cl_xhtml_aa5ab4cf85fe56de3821727cfcef327d9"><div class="ttname"><a href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a></div><div class="ttdeci">#define CONVOLUTION1x3_BIFROST2X1_STRIDE1(acc, src0, weights_row0)</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution_8cl_source.xhtml#l00211">depthwise_convolution.cl:211</a></div></div>
448<div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00142">helpers.h:142</a></div></div>
449<div class="ttc" id="helpers_8h_xhtml_a64d779f80eeb923e0ab2313433f7b40b"><div class="ttname"><a href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
450<div class="ttc" id="helpers_8h_xhtml_a7e4940407322d6f0ccb8b6b86b856019"><div class="ttname"><a href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a></div><div class="ttdeci">__global const uchar * vector_offset(const Vector *vec, int x)</div><div class="ttdoc">Get the pointer position of a Vector. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00290">helpers.h:290</a></div></div>
451<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
452<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="helpers_8h_source.xhtml#l00151">helpers.h:151</a></div></div>
453<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
454</div><!-- fragment -->
455</div>
456</div>
457<a class="anchor" id="a328ab548c536089e8238fb33de617bec"></a>
458<div class="memitem">
459<div class="memproto">
460 <table class="memname">
461 <tr>
462 <td class="memname">__kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost </td>
463 <td>(</td>
464 <td class="paramtype">__global uchar *&#160;</td>
465 <td class="paramname"><em>src_ptr</em>, </td>
466 </tr>
467 <tr>
468 <td class="paramkey"></td>
469 <td></td>
470 <td class="paramtype">uint&#160;</td>
471 <td class="paramname"><em>src_stride_x</em>, </td>
472 </tr>
473 <tr>
474 <td class="paramkey"></td>
475 <td></td>
476 <td class="paramtype">uint&#160;</td>
477 <td class="paramname"><em>src_step_x</em>, </td>
478 </tr>
479 <tr>
480 <td class="paramkey"></td>
481 <td></td>
482 <td class="paramtype">uint&#160;</td>
483 <td class="paramname"><em>src_stride_y</em>, </td>
484 </tr>
485 <tr>
486 <td class="paramkey"></td>
487 <td></td>
488 <td class="paramtype">uint&#160;</td>
489 <td class="paramname"><em>src_step_y</em>, </td>
490 </tr>
491 <tr>
492 <td class="paramkey"></td>
493 <td></td>
494 <td class="paramtype">uint&#160;</td>
495 <td class="paramname"><em>src_stride_z</em>, </td>
496 </tr>
497 <tr>
498 <td class="paramkey"></td>
499 <td></td>
500 <td class="paramtype">uint&#160;</td>
501 <td class="paramname"><em>src_step_z</em>, </td>
502 </tr>
503 <tr>
504 <td class="paramkey"></td>
505 <td></td>
506 <td class="paramtype">uint&#160;</td>
507 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
508 </tr>
509 <tr>
510 <td class="paramkey"></td>
511 <td></td>
512 <td class="paramtype">__global uchar *&#160;</td>
513 <td class="paramname"><em>dst_ptr</em>, </td>
514 </tr>
515 <tr>
516 <td class="paramkey"></td>
517 <td></td>
518 <td class="paramtype">uint&#160;</td>
519 <td class="paramname"><em>dst_stride_x</em>, </td>
520 </tr>
521 <tr>
522 <td class="paramkey"></td>
523 <td></td>
524 <td class="paramtype">uint&#160;</td>
525 <td class="paramname"><em>dst_step_x</em>, </td>
526 </tr>
527 <tr>
528 <td class="paramkey"></td>
529 <td></td>
530 <td class="paramtype">uint&#160;</td>
531 <td class="paramname"><em>dst_stride_y</em>, </td>
532 </tr>
533 <tr>
534 <td class="paramkey"></td>
535 <td></td>
536 <td class="paramtype">uint&#160;</td>
537 <td class="paramname"><em>dst_step_y</em>, </td>
538 </tr>
539 <tr>
540 <td class="paramkey"></td>
541 <td></td>
542 <td class="paramtype">uint&#160;</td>
543 <td class="paramname"><em>dst_stride_z</em>, </td>
544 </tr>
545 <tr>
546 <td class="paramkey"></td>
547 <td></td>
548 <td class="paramtype">uint&#160;</td>
549 <td class="paramname"><em>dst_step_z</em>, </td>
550 </tr>
551 <tr>
552 <td class="paramkey"></td>
553 <td></td>
554 <td class="paramtype">uint&#160;</td>
555 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
556 </tr>
557 <tr>
558 <td class="paramkey"></td>
559 <td></td>
560 <td class="paramtype">__global uchar *&#160;</td>
561 <td class="paramname"><em>weights_ptr</em>, </td>
562 </tr>
563 <tr>
564 <td class="paramkey"></td>
565 <td></td>
566 <td class="paramtype">uint&#160;</td>
567 <td class="paramname"><em>weights_stride_x</em>, </td>
568 </tr>
569 <tr>
570 <td class="paramkey"></td>
571 <td></td>
572 <td class="paramtype">uint&#160;</td>
573 <td class="paramname"><em>weights_step_x</em>, </td>
574 </tr>
575 <tr>
576 <td class="paramkey"></td>
577 <td></td>
578 <td class="paramtype">uint&#160;</td>
579 <td class="paramname"><em>weights_stride_y</em>, </td>
580 </tr>
581 <tr>
582 <td class="paramkey"></td>
583 <td></td>
584 <td class="paramtype">uint&#160;</td>
585 <td class="paramname"><em>weights_step_y</em>, </td>
586 </tr>
587 <tr>
588 <td class="paramkey"></td>
589 <td></td>
590 <td class="paramtype">uint&#160;</td>
591 <td class="paramname"><em>weights_stride_z</em>, </td>
592 </tr>
593 <tr>
594 <td class="paramkey"></td>
595 <td></td>
596 <td class="paramtype">uint&#160;</td>
597 <td class="paramname"><em>weights_step_z</em>, </td>
598 </tr>
599 <tr>
600 <td class="paramkey"></td>
601 <td></td>
602 <td class="paramtype">uint&#160;</td>
603 <td class="paramname"><em>weights_offset_first_element_in_bytes</em>, </td>
604 </tr>
605 <tr>
606 <td class="paramkey"></td>
607 <td></td>
608 <td class="paramtype">__global uchar *&#160;</td>
609 <td class="paramname"><em>biases_ptr</em>, </td>
610 </tr>
611 <tr>
612 <td class="paramkey"></td>
613 <td></td>
614 <td class="paramtype">uint&#160;</td>
615 <td class="paramname"><em>biases_stride_x</em>, </td>
616 </tr>
617 <tr>
618 <td class="paramkey"></td>
619 <td></td>
620 <td class="paramtype">uint&#160;</td>
621 <td class="paramname"><em>biases_step_x</em>, </td>
622 </tr>
623 <tr>
624 <td class="paramkey"></td>
625 <td></td>
626 <td class="paramtype">uint&#160;</td>
627 <td class="paramname"><em>biases_offset_first_element_in_bytes</em>&#160;</td>
628 </tr>
629 <tr>
630 <td></td>
631 <td>)</td>
632 <td></td><td></td>
633 </tr>
634 </table>
635</div><div class="memdoc">
636
637<p>This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 2. </p>
638<dl class="params"><dt>Parameters</dt><dd>
639 <table class="params">
640 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source image. Supported data types: F32 </td></tr>
641 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
642 <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>
643 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
644 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_y</td><td>src_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
645 <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 image </td></tr>
646 <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>
647 <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
648 <tr><td class="paramdir">[in]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F32 </td></tr>
649 <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>
650 <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>
651 <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>
652 <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>
653 <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>
654 <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>
655 <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>
656 <tr><td class="paramdir">[in]</td><td class="paramname">weights_ptr</td><td>Pointer to the weights tensor. Supported data types: F32 </td></tr>
657 <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_x</td><td>Stride of the weights tensor in X dimension (in bytes) </td></tr>
658 <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_x</td><td>weights_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
659 <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_y</td><td>Stride of the weights tensor in Y dimension (in bytes) </td></tr>
660 <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_y</td><td>weights_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
661 <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_z</td><td>Stride of the weights tensor in Z dimension (in bytes) </td></tr>
662 <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_z</td><td>weights_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
663 <tr><td class="paramdir">[in]</td><td class="paramname">weights_offset_first_element_in_bytes</td><td>The offset of the first element in the biases vector </td></tr>
664 <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>(Optional) Pointer to the biases vector. Supported data types: F32 </td></tr>
665 <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>(Optional) Stride of the biases vector in X dimension (in bytes) </td></tr>
666 <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>(Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
667 <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>(Optional) The offset of the first element in the biases vector </td></tr>
668 </table>
669 </dd>
670</dl>
671
672<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00360">360</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
673
674<p>References <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00117">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00096">CONVERT_TO_VECTOR_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00099">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>, <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00221">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00066">IMAGE_DECLARATION</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00134">Vector::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00151">Tensor3D::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00074">TENSOR3D_DECLARATION</a>, <a class="el" href="helpers_8h_source.xhtml#l00060">VECTOR_DECLARATION</a>, and <a class="el" href="helpers_8h_source.xhtml#l00290">vector_offset()</a>.</p>
675<div class="fragment"><div class="line"><a name="l00369"></a><span class="lineno"> 369</span>&#160;{</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00371"></a><span class="lineno"> 371</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="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> weights = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(weights);</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span>&#160;</div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span>&#160; float2 pixels0 = 0.0f;</div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>&#160; float2 pixels1 = 0.0f;</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>&#160;</div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>&#160; __global uchar *weights_addr = (__global uchar *)weights.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>;</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>&#160; __global uchar *src_addr = (__global uchar *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0);</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span>&#160;</div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>&#160; <span class="comment">// Load the weights</span></div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>&#160; float3 weights_row0 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 0 * weights_stride_y));</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>&#160; float3 weights_row1 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 1 * weights_stride_y));</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>&#160; float3 weights_row2 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 2 * weights_stride_y));</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span>&#160;</div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>&#160; <span class="comment">// Note: Since each work-item computes 4x2 elements, we need to load 5 rows from the input tensor</span></div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span>&#160; float4 src00 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span>&#160; float2 src01 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>&#160; float4 src10 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>&#160; float2 src11 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00390"></a><span class="lineno"> 390</span>&#160; float4 src20 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00391"></a><span class="lineno"> 391</span>&#160; float2 src21 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span>&#160; float4 src30 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span>&#160; float2 src31 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span>&#160; float4 src40 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row4</span></div><div class="line"><a name="l00395"></a><span class="lineno"> 395</span>&#160; float2 src41 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row4</span></div><div class="line"><a name="l00396"></a><span class="lineno"> 396</span>&#160;</div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src00, src01, weights_row0);</div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src10, src11, weights_row1);</div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src20, src21, weights_row2);</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src20, src21, weights_row0);</div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src30, src31, weights_row1);</div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span>&#160; <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src40, src41, weights_row2);</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>&#160;</div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>&#160;<span class="preprocessor">#ifdef HAS_BIAS</span></div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>&#160; <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>(biases);</div><div class="line"><a name="l00406"></a><span class="lineno"> 406</span>&#160;</div><div class="line"><a name="l00407"></a><span class="lineno"> 407</span>&#160; <span class="keywordtype">float</span> bias = *((__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a>(&amp;biases, get_global_id(2))));</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span>&#160;</div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>&#160; pixels0 += (float2)bias;</div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span>&#160; pixels1 += (float2)bias;</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/* defined(HAS_BIAS) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00412"></a><span class="lineno"> 412</span>&#160;</div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>&#160; vstore2(pixels0, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 0 * dst_stride_y));</div><div class="line"><a name="l00414"></a><span class="lineno"> 414</span>&#160; vstore2(pixels1, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 1 * dst_stride_y));</div><div class="line"><a name="l00415"></a><span class="lineno"> 415</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="helpers_8h_source.xhtml#l00132">helpers.h:132</a></div></div>
676<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#l00117">helpers.h:117</a></div></div>
677<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#l00149">helpers.h:149</a></div></div>
678<div class="ttc" id="depthwise__convolution_8cl_xhtml_a31cbdd6b28ee6331839432f69d87fdcb"><div class="ttname"><a href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a></div><div class="ttdeci">#define CONVOLUTION1x3_BIFROST2X1_STRIDE2(acc, src0, src1, weights_row0)</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution_8cl_source.xhtml#l00221">depthwise_convolution.cl:221</a></div></div>
679<div class="ttc" id="helpers_8h_xhtml_a009469e4d9b8fce3b6d5e97d2077827d"><div class="ttname"><a href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a></div><div class="ttdeci">__global uchar * offset(const Image *img, int x, int y)</div><div class="ttdoc">Get the pointer position of a Image. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00301">helpers.h:301</a></div></div>
680<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_convolution_8cpp_source.xhtml#l00137">Convolution.cpp:137</a></div></div>
681<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="helpers_8h_source.xhtml#l00140">helpers.h:140</a></div></div>
682<div class="ttc" id="struct_image_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">Image::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00142">helpers.h:142</a></div></div>
683<div class="ttc" id="helpers_8h_xhtml_a64d779f80eeb923e0ab2313433f7b40b"><div class="ttname"><a href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
684<div class="ttc" id="helpers_8h_xhtml_a7e4940407322d6f0ccb8b6b86b856019"><div class="ttname"><a href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a></div><div class="ttdeci">__global const uchar * vector_offset(const Vector *vec, int x)</div><div class="ttdoc">Get the pointer position of a Vector. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00290">helpers.h:290</a></div></div>
685<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
686<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="helpers_8h_source.xhtml#l00151">helpers.h:151</a></div></div>
687<div class="ttc" id="namespacearm__compute_1_1test_1_1validation_xhtml_a6743f0a130e8311e6f5b1a23df102472"><div class="ttname"><a href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">arm_compute::test::validation::src</a></div><div class="ttdeci">convolution configure &amp; src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
688</div><!-- fragment -->
689</div>
690</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]691</div><!-- contents -->
692</div><!-- doc-content -->
693<!-- start footer part -->
694<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
695 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]696 <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="depthwise__convolution_8cl.xhtml">depthwise_convolution.cl</a></li>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000[diff] [blame^]697 <li class="footer">Generated on Thu Feb 22 2018 15:45:22 for Compute Library by
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]698 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame]699 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]700 </ul>
701</div>
702</body>
703</html>