blob: aacc49c6ae46db6ba4bba8f02680040d63800fc4 [file] [log] [blame]
Anthony Barbierdbdab852017-06-23 15:42:00 +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"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +00007<meta name="generator" content="Doxygen 1.8.11"/>
Anthony Barbierdbdab852017-06-23 15:42:00 +01008<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
9<title>Compute Library: src/core/CL/cl_kernels/concatenate.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>
Anthony Barbierdbdab852017-06-23 15:42:00 +010016<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 +000022<script type="text/javascript" src="search/searchdata.js"></script>
Anthony Barbierdbdab852017-06-23 15:42:00 +010023<script type="text/javascript" src="search/search.js"></script>
24<script type="text/javascript">
Anthony Barbier8140e1e2017-12-14 23:48:46 +000025 $(document).ready(function() { init_search(); });
Anthony Barbierdbdab852017-06-23 15:42:00 +010026</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 +000032</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Anthony Barbierdbdab852017-06-23 15:42:00 +010033<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
Jenkinsb3a371b2018-05-23 11:36:53 +010043 &#160;<span id="projectnumber">18.05</span>
Anthony Barbierdbdab852017-06-23 15:42:00 +010044 </div>
45 </td>
46 </tr>
47 </tbody>
48</table>
49</div>
50<!-- end header part -->
Anthony Barbier8140e1e2017-12-14 23:48:46 +000051<!-- Generated by Doxygen 1.8.11 -->
Anthony Barbierdbdab852017-06-23 15:42:00 +010052<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('concatenate_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 +0000106</div>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100107
108<!-- iframe showing the search results (closed by default) -->
109<div id="MSearchResultsWindow">
110<iframe src="javascript:void(0)" frameborder="0"
111 name="MSearchResults" id="MSearchResults">
112</iframe>
113</div>
114
115<div class="header">
116 <div class="summary">
117<a href="#func-members">Functions</a> </div>
118 <div class="headertitle">
119<div class="title">concatenate.cl File Reference</div> </div>
120</div><!--header-->
121<div class="contents">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000122<div class="textblock"><code>#include &quot;<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
Anthony Barbierdbdab852017-06-23 15:42:00 +0100123</div>
124<p><a href="concatenate_8cl_source.xhtml">Go to the source code of this file.</a></p>
125<table class="memberdecls">
126<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
127Functions</h2></td></tr>
Jenkinsb3a371b2018-05-23 11:36:53 +0100128<tr class="memitem:ae99c297ac202e9a2fc1c622925dea12e"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="concatenate_8cl.xhtml#ae99c297ac202e9a2fc1c622925dea12e">concatenate_width</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, int <a class="el" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>)</td></tr>
129<tr class="memdesc:ae99c297ac202e9a2fc1c622925dea12e"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel concatenates the input tensor into the output tensor along the first dimension. <a href="#ae99c297ac202e9a2fc1c622925dea12e">More...</a><br /></td></tr>
130<tr class="separator:ae99c297ac202e9a2fc1c622925dea12e"><td class="memSeparator" colspan="2">&#160;</td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100131<tr class="memitem:a790788f3dd634f6ca93e2f9ffeaa0ed5"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="concatenate_8cl.xhtml#a790788f3dd634f6ca93e2f9ffeaa0ed5">concatenate_depth</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, int3 offsets)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000132<tr class="memdesc:a790788f3dd634f6ca93e2f9ffeaa0ed5"><td class="mdescLeft">&#160;</td><td class="mdescRight">This kernel concatenates the input tensor into the output tensor along the third dimension. <a href="#a790788f3dd634f6ca93e2f9ffeaa0ed5">More...</a><br /></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100133<tr class="separator:a790788f3dd634f6ca93e2f9ffeaa0ed5"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100134</table>
135<h2 class="groupheader">Function Documentation</h2>
Kaizen8938bd32017-09-28 14:38:23 +0100136<a class="anchor" id="a790788f3dd634f6ca93e2f9ffeaa0ed5"></a>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100137<div class="memitem">
138<div class="memproto">
139 <table class="memname">
140 <tr>
141 <td class="memname">__kernel void concatenate_depth </td>
142 <td>(</td>
143 <td class="paramtype">__global uchar *&#160;</td>
144 <td class="paramname"><em>src_ptr</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_stride_x</em>, </td>
151 </tr>
152 <tr>
153 <td class="paramkey"></td>
154 <td></td>
155 <td class="paramtype">uint&#160;</td>
156 <td class="paramname"><em>src_step_x</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>src_stride_y</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>src_step_y</em>, </td>
169 </tr>
170 <tr>
171 <td class="paramkey"></td>
172 <td></td>
173 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100174 <td class="paramname"><em>src_stride_z</em>, </td>
175 </tr>
176 <tr>
177 <td class="paramkey"></td>
178 <td></td>
179 <td class="paramtype">uint&#160;</td>
180 <td class="paramname"><em>src_step_z</em>, </td>
181 </tr>
182 <tr>
183 <td class="paramkey"></td>
184 <td></td>
185 <td class="paramtype">uint&#160;</td>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100186 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
187 </tr>
188 <tr>
189 <td class="paramkey"></td>
190 <td></td>
191 <td class="paramtype">__global uchar *&#160;</td>
192 <td class="paramname"><em>dst_ptr</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_stride_x</em>, </td>
199 </tr>
200 <tr>
201 <td class="paramkey"></td>
202 <td></td>
203 <td class="paramtype">uint&#160;</td>
204 <td class="paramname"><em>dst_step_x</em>, </td>
205 </tr>
206 <tr>
207 <td class="paramkey"></td>
208 <td></td>
209 <td class="paramtype">uint&#160;</td>
210 <td class="paramname"><em>dst_stride_y</em>, </td>
211 </tr>
212 <tr>
213 <td class="paramkey"></td>
214 <td></td>
215 <td class="paramtype">uint&#160;</td>
216 <td class="paramname"><em>dst_step_y</em>, </td>
217 </tr>
218 <tr>
219 <td class="paramkey"></td>
220 <td></td>
221 <td class="paramtype">uint&#160;</td>
Kaizen8938bd32017-09-28 14:38:23 +0100222 <td class="paramname"><em>dst_stride_z</em>, </td>
223 </tr>
224 <tr>
225 <td class="paramkey"></td>
226 <td></td>
227 <td class="paramtype">uint&#160;</td>
228 <td class="paramname"><em>dst_step_z</em>, </td>
229 </tr>
230 <tr>
231 <td class="paramkey"></td>
232 <td></td>
233 <td class="paramtype">uint&#160;</td>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100234 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
235 </tr>
236 <tr>
237 <td class="paramkey"></td>
238 <td></td>
Kaizen8938bd32017-09-28 14:38:23 +0100239 <td class="paramtype">int3&#160;</td>
240 <td class="paramname"><em>offsets</em>&#160;</td>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100241 </tr>
242 <tr>
243 <td></td>
244 <td>)</td>
245 <td></td><td></td>
246 </tr>
247 </table>
248</div><div class="memdoc">
249
250<p>This kernel concatenates the input tensor into the output tensor along the third dimension. </p>
251<dl class="params"><dt>Parameters</dt><dd>
252 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100253 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source tensor. Supported data types: QS8, QS16, F16, F32 </td></tr>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100254 <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>
255 <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>
256 <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>
257 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100258 <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>
259 <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>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100260 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100261 <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>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100262 <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>
263 <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>
264 <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>
265 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100266 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
267 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Z processed per workitem(in bytes) </td></tr>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100268 <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>
Kaizen8938bd32017-09-28 14:38:23 +0100269 <tr><td class="paramdir">[in]</td><td class="paramname">offsets</td><td>The offsets to the first valid element of the output tensor in bytes </td></tr>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100270 </table>
271 </dd>
272</dl>
273
Jenkinsb3a371b2018-05-23 11:36:53 +0100274<p>Definition at line <a class="el" href="concatenate_8cl_source.xhtml#l00081">81</a> of file <a class="el" href="concatenate_8cl_source.xhtml">concatenate.cl</a>.</p>
275
276<p>References <a class="el" href="helpers_8h_source.xhtml#l00119">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00137">arm_compute::test::validation::dst</a>, <a class="el" href="helpers_8h_source.xhtml#l00153">Tensor3D::ptr</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00147">arm_compute::test::validation::src</a>, <a class="el" href="helpers_8h_source.xhtml#l00315">tensor3D_offset()</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00042">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00045">VSTORE</a>.</p>
277<div class="fragment"><div class="line"><a name="l00085"></a><span class="lineno"> 85</span>&#160;{</div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span>&#160; <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(src);</div><div class="line"><a name="l00087"></a><span class="lineno"> 87</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="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(dst);</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span>&#160;</div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span>&#160; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, VEC_SIZE)</div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>&#160; source_values = <a class="code" href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(VEC_SIZE)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)<a class="code" href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a>(&amp;src, -offsets.x, -offsets.y, 0));</div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span>&#160;</div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VEC_SIZE)</div><div class="line"><a name="l00093"></a><span class="lineno"> 93</span>&#160; (source_values, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)(dst.ptr + offsets.z));</div><div class="line"><a name="l00094"></a><span class="lineno"> 94</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_a287e2fc366c312b468382c95bb90f91f"><div class="ttname"><a href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a></div><div class="ttdeci">#define VLOAD(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00042">helpers.h:42</a></div></div>
278<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#l00119">helpers.h:119</a></div></div>
279<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>
280<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#l00151">helpers.h:151</a></div></div>
281<div class="ttc" id="helpers_8h_xhtml_acb282042d1edeeaa3cc979a206f78b54"><div class="ttname"><a href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a></div><div class="ttdeci">#define VSTORE(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00045">helpers.h:45</a></div></div>
282<div class="ttc" id="helpers_8h_xhtml_a2101b2fe0193ce227ae4e0945e321d85"><div class="ttname"><a href="helpers_8h.xhtml#a2101b2fe0193ce227ae4e0945e321d85">tensor3D_offset</a></div><div class="ttdeci">__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)</div><div class="ttdoc">Get the pointer position of a Tensor3D. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00315">helpers.h:315</a></div></div>
283<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>
284<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_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="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
285<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>
286</div><!-- fragment -->
287</div>
288</div>
289<a class="anchor" id="ae99c297ac202e9a2fc1c622925dea12e"></a>
290<div class="memitem">
291<div class="memproto">
292 <table class="memname">
293 <tr>
294 <td class="memname">__kernel void concatenate_width </td>
295 <td>(</td>
296 <td class="paramtype">__global uchar *&#160;</td>
297 <td class="paramname"><em>src_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>src_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>src_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>src_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>src_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>src_stride_z</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>src_step_z</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>src_offset_first_element_in_bytes</em>, </td>
340 </tr>
341 <tr>
342 <td class="paramkey"></td>
343 <td></td>
344 <td class="paramtype">__global uchar *&#160;</td>
345 <td class="paramname"><em>dst_ptr</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_x</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_x</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_stride_y</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>dst_step_y</em>, </td>
370 </tr>
371 <tr>
372 <td class="paramkey"></td>
373 <td></td>
374 <td class="paramtype">uint&#160;</td>
375 <td class="paramname"><em>dst_stride_z</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>dst_step_z</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>dst_offset_first_element_in_bytes</em>, </td>
388 </tr>
389 <tr>
390 <td class="paramkey"></td>
391 <td></td>
392 <td class="paramtype">int&#160;</td>
393 <td class="paramname"><em>offset</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 kernel concatenates the input tensor into the output tensor along the first dimension. </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 tensor. Supported data types: QS8, QASYMM8, QS16, F16, F32 </td></tr>
407 <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>
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 tensor 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_stride_z</td><td>Stride of the source tensor in Z dimension (in bytes) </td></tr>
412 <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>
413 <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>
414 <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>
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 source 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 Z 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">offset</td><td>The offset to the first valid element of the output tensor in bytes </td></tr>
423 </table>
424 </dd>
425</dl>
426
Kaizen8938bd32017-09-28 14:38:23 +0100427<p>Definition at line <a class="el" href="concatenate_8cl_source.xhtml#l00046">46</a> of file <a class="el" href="concatenate_8cl_source.xhtml">concatenate.cl</a>.</p>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100428
Jenkinsb3a371b2018-05-23 11:36:53 +0100429<p>References <a class="el" href="helpers_8h_source.xhtml#l00119">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00137">arm_compute::test::validation::dst</a>, <a class="el" href="helpers_8h_source.xhtml#l00153">Tensor3D::ptr</a>, <a class="el" href="_c_l_2_convolution_8cpp_source.xhtml#l00147">arm_compute::test::validation::src</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00042">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00045">VSTORE</a>.</p>
430<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_tensor3_d.xhtml">Tensor3D</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(src);</div><div class="line"><a name="l00052"></a><span class="lineno"> 52</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="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(dst);</div><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; <a class="code" href="fixed__point_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a>, VEC_SIZE)</div><div class="line"><a name="l00055"></a><span class="lineno"> 55</span>&#160; source_values = <a class="code" href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a>(VEC_SIZE)(0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)src.ptr);</div><div class="line"><a name="l00056"></a><span class="lineno"> 56</span>&#160;</div><div class="line"><a name="l00057"></a><span class="lineno"> 57</span>&#160; <a class="code" href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a>(VEC_SIZE)</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>&#160; (source_values, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#afb8c72ce35c4a1f4a2588d6573e54aa1">DATA_TYPE</a> *)(dst.ptr + <a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>));</div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_a287e2fc366c312b468382c95bb90f91f"><div class="ttname"><a href="helpers_8h.xhtml#a287e2fc366c312b468382c95bb90f91f">VLOAD</a></div><div class="ttdeci">#define VLOAD(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00042">helpers.h:42</a></div></div>
431<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#l00119">helpers.h:119</a></div></div>
Kaizen8938bd32017-09-28 14:38:23 +0100432<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>
Jenkinsb3a371b2018-05-23 11:36:53 +0100433<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#l00151">helpers.h:151</a></div></div>
434<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#l00303">helpers.h:303</a></div></div>
435<div class="ttc" id="helpers_8h_xhtml_acb282042d1edeeaa3cc979a206f78b54"><div class="ttname"><a href="helpers_8h.xhtml#acb282042d1edeeaa3cc979a206f78b54">VSTORE</a></div><div class="ttdeci">#define VSTORE(size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00045">helpers.h:45</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000436<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>
Kaizen8938bd32017-09-28 14:38:23 +0100437<div class="ttc" id="fixed__point_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="fixed__point_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="fixed__point_8h_source.xhtml#l00093">fixed_point.h:93</a></div></div>
Anthony Barbier06ea0482018-02-22 15:45:35 +0000438<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>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100439</div><!-- fragment -->
440</div>
441</div>
442</div><!-- contents -->
443</div><!-- doc-content -->
444<!-- start footer part -->
445<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
446 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000447 <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="concatenate_8cl.xhtml">concatenate.cl</a></li>
Jenkinsb3a371b2018-05-23 11:36:53 +0100448 <li class="footer">Generated on Wed May 23 2018 11:36:38 for Compute Library by
Anthony Barbierdbdab852017-06-23 15:42:00 +0100449 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000450 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100451 </ul>
452</div>
453</body>
454</html>