Gitiles
Code Review Sign In
gerrit-public.fairphone.software / platform / external / ComputeLibrary / 8140e1e155d3430992fa46e04ef8938ff09ffd2d / . / documentation / gemm_8cl.xhtml
blob: 0f56ac00eaaac0b1de2377dfa711585e02789ca9 [file] [log] [blame]
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]1<!-- HTML header for doxygen 1.8.9.1-->
2<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
3<html xmlns="http://www.w3.org/1999/xhtml">
4<head>
5<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
6<meta http-equiv="X-UA-Compatible" content="IE=9"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]7<meta name="generator" content="Doxygen 1.8.11"/>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]8<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]9<title>Compute Library: src/core/CL/cl_kernels/gemm.cl File Reference</title>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]10<link href="tabs.css" rel="stylesheet" type="text/css"/>
11<script type="text/javascript" src="jquery.js"></script>
12<script type="text/javascript" src="dynsections.js"></script>
13<link href="navtree.css" rel="stylesheet" type="text/css"/>
14<script type="text/javascript" src="resize.js"></script>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]15<script type="text/javascript" src="navtreedata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]16<script type="text/javascript" src="navtree.js"></script>
17<script type="text/javascript">
18 $(document).ready(initResizable);
19 $(window).load(resizeHeight);
20</script>
21<link href="search/search.css" rel="stylesheet" type="text/css"/>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]22<script type="text/javascript" src="search/searchdata.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]23<script type="text/javascript" src="search/search.js"></script>
24<script type="text/javascript">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]25 $(document).ready(function() { init_search(); });
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]26</script>
27<script type="text/x-mathjax-config">
28 MathJax.Hub.Config({
29 extensions: ["tex2jax.js"],
30 jax: ["input/TeX","output/HTML-CSS"],
31});
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]32</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]33<link href="doxygen.css" rel="stylesheet" type="text/css" />
34</head>
35<body>
36<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
37<div id="titlearea">
38<table cellspacing="0" cellpadding="0">
39 <tbody>
40 <tr style="height: 56px;">
41 <td style="padding-left: 0.5em;">
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]42 <div id="projectname">Compute Library
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]43 &#160;<span id="projectnumber">17.12</span>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]44 </div>
45 </td>
46 </tr>
47 </tbody>
48</table>
49</div>
50<!-- end header part -->
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]51<!-- Generated by Doxygen 1.8.11 -->
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]52<script type="text/javascript">
53var searchBox = new SearchBox("searchBox", "search",false,'Search');
54</script>
55 <div id="navrow1" class="tabs">
56 <ul class="tablist">
57 <li><a href="index.xhtml"><span>Main&#160;Page</span></a></li>
Anthony Barbierdbdab852017-06-23 15:42:00 +0100[diff] [blame]58 <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]59 <li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
60 <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
61 <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
62 <li>
63 <div id="MSearchBox" class="MSearchBoxInactive">
64 <span class="left">
65 <img id="MSearchSelect" src="search/mag_sel.png"
66 onmouseover="return searchBox.OnSearchSelectShow()"
67 onmouseout="return searchBox.OnSearchSelectHide()"
68 alt=""/>
69 <input type="text" id="MSearchField" value="Search" accesskey="S"
70 onfocus="searchBox.OnSearchFieldFocus(true)"
71 onblur="searchBox.OnSearchFieldFocus(false)"
72 onkeyup="searchBox.OnSearchFieldChange(event)"/>
73 </span><span class="right">
74 <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
75 </span>
76 </div>
77 </li>
78 </ul>
79 </div>
80 <div id="navrow2" class="tabs2">
81 <ul class="tablist">
82 <li><a href="files.xhtml"><span>File&#160;List</span></a></li>
83 <li><a href="globals.xhtml"><span>Globals</span></a></li>
84 </ul>
85 </div>
86</div><!-- top -->
87<div id="side-nav" class="ui-resizable side-nav-resizable">
88 <div id="nav-tree">
89 <div id="nav-tree-contents">
90 <div id="nav-sync" class="sync"></div>
91 </div>
92 </div>
93 <div id="splitbar" style="-moz-user-select:none;"
94 class="ui-resizable-handle">
95 </div>
96</div>
97<script type="text/javascript">
98$(document).ready(function(){initNavTree('gemm_8cl.xhtml','');});
99</script>
100<div id="doc-content">
101<!-- window showing the filter options -->
102<div id="MSearchSelectWindow"
103 onmouseover="return searchBox.OnSearchSelectShow()"
104 onmouseout="return searchBox.OnSearchSelectHide()"
105 onkeydown="return searchBox.OnSearchSelectKey(event)">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]106</div>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]107
108<!-- iframe showing the search results (closed by default) -->
109<div id="MSearchResultsWindow">
110<iframe src="javascript:void(0)" frameborder="0"
111 name="MSearchResults" id="MSearchResults">
112</iframe>
113</div>
114
115<div class="header">
116 <div class="summary">
117<a href="#func-members">Functions</a> </div>
118 <div class="headertitle">
119<div class="title">gemm.cl File Reference</div> </div>
120</div><!--header-->
121<div class="contents">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]122<div class="textblock"><code>#include &quot;<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>&quot;</code><br />
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]123</div>
124<p><a href="gemm_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>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]128<tr class="memitem:ac5eba21b9ac015e1454044aeba93546c"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#ac5eba21b9ac015e1454044aeba93546c">gemm_transpose1x4</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]129<tr class="memdesc:ac5eba21b9ac015e1454044aeba93546c"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel computes the "vector" 1x4 transposition of input matrix. <a href="#ac5eba21b9ac015e1454044aeba93546c">More...</a><br /></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]130<tr class="separator:ac5eba21b9ac015e1454044aeba93546c"><td class="memSeparator" colspan="2">&#160;</td></tr>
131<tr class="memitem:a75c35ac120f123e9cce56c43ba3f48d2"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a75c35ac120f123e9cce56c43ba3f48d2">gemm_transpose1x8</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]132<tr class="memdesc:a75c35ac120f123e9cce56c43ba3f48d2"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel computes the "vector" 1x8 transposition of input matrix. <a href="#a75c35ac120f123e9cce56c43ba3f48d2">More...</a><br /></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]133<tr class="separator:a75c35ac120f123e9cce56c43ba3f48d2"><td class="memSeparator" colspan="2">&#160;</td></tr>
134<tr class="memitem:a97a299ffb620aa98fe5e39bb8ab0833f"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a97a299ffb620aa98fe5e39bb8ab0833f">gemm_transpose1x16</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]135<tr class="memdesc:a97a299ffb620aa98fe5e39bb8ab0833f"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel computes the "vector" 1x16 transposition of input matrix. <a href="#a97a299ffb620aa98fe5e39bb8ab0833f">More...</a><br /></td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]136<tr class="separator:a97a299ffb620aa98fe5e39bb8ab0833f"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]137<tr class="memitem:a560b2d50aa886edc0f5daf4fe729717f"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a560b2d50aa886edc0f5daf4fe729717f">gemm_interleave4x4_32bit</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]138<tr class="memdesc:a560b2d50aa886edc0f5daf4fe729717f"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. <a href="#a560b2d50aa886edc0f5daf4fe729717f">More...</a><br /></td></tr>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]139<tr class="separator:a560b2d50aa886edc0f5daf4fe729717f"><td class="memSeparator" colspan="2">&#160;</td></tr>
140<tr class="memitem:a2f32d740e780059f68da8aa589ed0a5b"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#a2f32d740e780059f68da8aa589ed0a5b">gemm_interleave4x4_16bit</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]141<tr class="memdesc:a2f32d740e780059f68da8aa589ed0a5b"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. <a href="#a2f32d740e780059f68da8aa589ed0a5b">More...</a><br /></td></tr>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]142<tr class="separator:a2f32d740e780059f68da8aa589ed0a5b"><td class="memSeparator" colspan="2">&#160;</td></tr>
143<tr class="memitem:aa66a7b6b0420e54ec173743e6c5bfc45"><td class="memItemLeft" align="right" valign="top">__kernel void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="gemm_8cl.xhtml#aa66a7b6b0420e54ec173743e6c5bfc45">gemm_interleave4x4_8bit</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)</td></tr>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]144<tr class="memdesc:aa66a7b6b0420e54ec173743e6c5bfc45"><td class="mdescLeft">&#160;</td><td class="mdescRight">This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. <a href="#aa66a7b6b0420e54ec173743e6c5bfc45">More...</a><br /></td></tr>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]145<tr class="separator:aa66a7b6b0420e54ec173743e6c5bfc45"><td class="memSeparator" colspan="2">&#160;</td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]146</table>
147<h2 class="groupheader">Function Documentation</h2>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]148<a class="anchor" id="a2f32d740e780059f68da8aa589ed0a5b"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]149<div class="memitem">
150<div class="memproto">
151 <table class="memname">
152 <tr>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]153 <td class="memname">__kernel void gemm_interleave4x4_16bit </td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]154 <td>(</td>
155 <td class="paramtype">__global uchar *&#160;</td>
156 <td class="paramname"><em>src_ptr</em>, </td>
157 </tr>
158 <tr>
159 <td class="paramkey"></td>
160 <td></td>
161 <td class="paramtype">uint&#160;</td>
162 <td class="paramname"><em>src_stride_x</em>, </td>
163 </tr>
164 <tr>
165 <td class="paramkey"></td>
166 <td></td>
167 <td class="paramtype">uint&#160;</td>
168 <td class="paramname"><em>src_step_x</em>, </td>
169 </tr>
170 <tr>
171 <td class="paramkey"></td>
172 <td></td>
173 <td class="paramtype">uint&#160;</td>
174 <td class="paramname"><em>src_stride_y</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_y</em>, </td>
181 </tr>
182 <tr>
183 <td class="paramkey"></td>
184 <td></td>
185 <td class="paramtype">uint&#160;</td>
186 <td class="paramname"><em>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>
222 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
223 </tr>
224 <tr>
225 <td></td>
226 <td>)</td>
227 <td></td><td></td>
228 </tr>
229 </table>
230</div><div class="memdoc">
231
232<p>This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. </p>
233<dl class="params"><dt>Parameters</dt><dd>
234 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]235 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U16/S16/QS16/F16 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]236 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
237 <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>
238 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
239 <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>
240 <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 matrix </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]241 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]242 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
243 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
244 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
245 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
246 <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 matrix </td></tr>
247 </table>
248 </dd>
249</dl>
250
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]251<p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00182">182</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]252
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]253<p>References <a class="el" href="helpers_8h_source.xhtml#l00102">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">arm_compute::test::validation::dst</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>.</p>
254<div class="fragment"><div class="line"><a name="l00184"></a><span class="lineno"> 184</span>&#160;{</div><div class="line"><a name="l00185"></a><span class="lineno"> 185</span>&#160; <span class="comment">// Compute source and destination addresses</span></div><div class="line"><a name="l00186"></a><span class="lineno"> 186</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00187"></a><span class="lineno"> 187</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#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>&#160;</div><div class="line"><a name="l00189"></a><span class="lineno"> 189</span>&#160; <span class="comment">// Load values from Matrix A</span></div><div class="line"><a name="l00190"></a><span class="lineno"> 190</span>&#160; ushort8 a0 = vload8(0, (__global ushort *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0)));</div><div class="line"><a name="l00191"></a><span class="lineno"> 191</span>&#160; ushort8 a1 = vload8(0, (__global ushort *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 1)));</div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>&#160; ushort8 a2 = vload8(0, (__global ushort *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 2)));</div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span>&#160; ushort8 a3 = vload8(0, (__global ushort *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 3)));</div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>&#160;</div><div class="line"><a name="l00195"></a><span class="lineno"> 195</span>&#160; ushort8 val0 = (ushort8)((ushort4)(a0.s0, a1.s0, a2.s0, a3.s0), (ushort4)(a0.s1, a1.s1, a2.s1, a3.s1));</div><div class="line"><a name="l00196"></a><span class="lineno"> 196</span>&#160; vstore8(val0, 0, ((__global ushort *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 0);</div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>&#160;</div><div class="line"><a name="l00198"></a><span class="lineno"> 198</span>&#160; val0 = (ushort8)((ushort4)(a0.s2, a1.s2, a2.s2, a3.s2), (ushort4)(a0.s3, a1.s3, a2.s3, a3.s3));</div><div class="line"><a name="l00199"></a><span class="lineno"> 199</span>&#160; vstore8(val0, 0, ((__global ushort *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 8);</div><div class="line"><a name="l00200"></a><span class="lineno"> 200</span>&#160;</div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span>&#160; val0 = (ushort8)((ushort4)(a0.s4, a1.s4, a2.s4, a3.s4), (ushort4)(a0.s5, a1.s5, a2.s5, a3.s5));</div><div class="line"><a name="l00202"></a><span class="lineno"> 202</span>&#160; vstore8(val0, 0, ((__global ushort *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 16);</div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>&#160;</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</span>&#160; val0 = (ushort8)((ushort4)(a0.s6, a1.s6, a2.s6, a3.s6), (ushort4)(a0.s7, a1.s7, a2.s7, a3.s7));</div><div class="line"><a name="l00205"></a><span class="lineno"> 205</span>&#160; vstore8(val0, 0, ((__global ushort *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 24);</div><div class="line"><a name="l00206"></a><span class="lineno"> 206</span>&#160;}</div><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>
255<div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00102">helpers.h:102</a></div></div>
256<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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
257<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>
258<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]259</div><!-- fragment -->
260</div>
261</div>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]262<a class="anchor" id="a560b2d50aa886edc0f5daf4fe729717f"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]263<div class="memitem">
264<div class="memproto">
265 <table class="memname">
266 <tr>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]267 <td class="memname">__kernel void gemm_interleave4x4_32bit </td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]268 <td>(</td>
269 <td class="paramtype">__global uchar *&#160;</td>
270 <td class="paramname"><em>src_ptr</em>, </td>
271 </tr>
272 <tr>
273 <td class="paramkey"></td>
274 <td></td>
275 <td class="paramtype">uint&#160;</td>
276 <td class="paramname"><em>src_stride_x</em>, </td>
277 </tr>
278 <tr>
279 <td class="paramkey"></td>
280 <td></td>
281 <td class="paramtype">uint&#160;</td>
282 <td class="paramname"><em>src_step_x</em>, </td>
283 </tr>
284 <tr>
285 <td class="paramkey"></td>
286 <td></td>
287 <td class="paramtype">uint&#160;</td>
288 <td class="paramname"><em>src_stride_y</em>, </td>
289 </tr>
290 <tr>
291 <td class="paramkey"></td>
292 <td></td>
293 <td class="paramtype">uint&#160;</td>
294 <td class="paramname"><em>src_step_y</em>, </td>
295 </tr>
296 <tr>
297 <td class="paramkey"></td>
298 <td></td>
299 <td class="paramtype">uint&#160;</td>
300 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
301 </tr>
302 <tr>
303 <td class="paramkey"></td>
304 <td></td>
305 <td class="paramtype">__global uchar *&#160;</td>
306 <td class="paramname"><em>dst_ptr</em>, </td>
307 </tr>
308 <tr>
309 <td class="paramkey"></td>
310 <td></td>
311 <td class="paramtype">uint&#160;</td>
312 <td class="paramname"><em>dst_stride_x</em>, </td>
313 </tr>
314 <tr>
315 <td class="paramkey"></td>
316 <td></td>
317 <td class="paramtype">uint&#160;</td>
318 <td class="paramname"><em>dst_step_x</em>, </td>
319 </tr>
320 <tr>
321 <td class="paramkey"></td>
322 <td></td>
323 <td class="paramtype">uint&#160;</td>
324 <td class="paramname"><em>dst_stride_y</em>, </td>
325 </tr>
326 <tr>
327 <td class="paramkey"></td>
328 <td></td>
329 <td class="paramtype">uint&#160;</td>
330 <td class="paramname"><em>dst_step_y</em>, </td>
331 </tr>
332 <tr>
333 <td class="paramkey"></td>
334 <td></td>
335 <td class="paramtype">uint&#160;</td>
336 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
337 </tr>
338 <tr>
339 <td></td>
340 <td>)</td>
341 <td></td><td></td>
342 </tr>
343 </table>
344</div><div class="memdoc">
345
346<p>This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. </p>
347<dl class="params"><dt>Parameters</dt><dd>
348 <table class="params">
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]349 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U32/S32/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]350 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
351 <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>
352 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
353 <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>
354 <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 matrix </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]355 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]356 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
357 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
358 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
359 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
360 <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 matrix </td></tr>
361 </table>
362 </dd>
363</dl>
364
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]365<p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00141">141</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]366
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]367<p>References <a class="el" href="helpers_8h_source.xhtml#l00102">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">arm_compute::test::validation::dst</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, and <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>.</p>
368<div class="fragment"><div class="line"><a name="l00143"></a><span class="lineno"> 143</span>&#160;{</div><div class="line"><a name="l00144"></a><span class="lineno"> 144</span>&#160; <span class="comment">// Compute source and destination addresses</span></div><div class="line"><a name="l00145"></a><span class="lineno"> 145</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00146"></a><span class="lineno"> 146</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#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00147"></a><span class="lineno"> 147</span>&#160;</div><div class="line"><a name="l00148"></a><span class="lineno"> 148</span>&#160; <span class="comment">// Load values from Matrix A</span></div><div class="line"><a name="l00149"></a><span class="lineno"> 149</span>&#160; uint4 a0 = vload4(0, (__global uint *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0)));</div><div class="line"><a name="l00150"></a><span class="lineno"> 150</span>&#160; uint4 a1 = vload4(0, (__global uint *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 1)));</div><div class="line"><a name="l00151"></a><span class="lineno"> 151</span>&#160; uint4 a2 = vload4(0, (__global uint *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 2)));</div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span>&#160; uint4 a3 = vload4(0, (__global uint *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 3)));</div><div class="line"><a name="l00153"></a><span class="lineno"> 153</span>&#160;</div><div class="line"><a name="l00154"></a><span class="lineno"> 154</span>&#160; uint4 val0 = (uint4)(a0.s0, a1.s0, a2.s0, a3.s0);</div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span>&#160; vstore4(val0, 0, ((__global uint *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 0);</div><div class="line"><a name="l00156"></a><span class="lineno"> 156</span>&#160;</div><div class="line"><a name="l00157"></a><span class="lineno"> 157</span>&#160; val0 = (uint4)(a0.s1, a1.s1, a2.s1, a3.s1);</div><div class="line"><a name="l00158"></a><span class="lineno"> 158</span>&#160; vstore4(val0, 0, ((__global uint *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 4);</div><div class="line"><a name="l00159"></a><span class="lineno"> 159</span>&#160;</div><div class="line"><a name="l00160"></a><span class="lineno"> 160</span>&#160; val0 = (uint4)(a0.s2, a1.s2, a2.s2, a3.s2);</div><div class="line"><a name="l00161"></a><span class="lineno"> 161</span>&#160; vstore4(val0, 0, ((__global uint *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 8);</div><div class="line"><a name="l00162"></a><span class="lineno"> 162</span>&#160;</div><div class="line"><a name="l00163"></a><span class="lineno"> 163</span>&#160; val0 = (uint4)(a0.s3, a1.s3, a2.s3, a3.s3);</div><div class="line"><a name="l00164"></a><span class="lineno"> 164</span>&#160; vstore4(val0, 0, ((__global uint *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 12);</div><div class="line"><a name="l00165"></a><span class="lineno"> 165</span>&#160;}</div><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>
369<div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00102">helpers.h:102</a></div></div>
370<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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
371<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>
372<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]373</div><!-- fragment -->
374</div>
375</div>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]376<a class="anchor" id="aa66a7b6b0420e54ec173743e6c5bfc45"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]377<div class="memitem">
378<div class="memproto">
379 <table class="memname">
380 <tr>
Anthony Barbier46d59272017-05-04 09:15:15 +0100[diff] [blame]381 <td class="memname">__kernel void gemm_interleave4x4_8bit </td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]382 <td>(</td>
383 <td class="paramtype">__global uchar *&#160;</td>
384 <td class="paramname"><em>src_ptr</em>, </td>
385 </tr>
386 <tr>
387 <td class="paramkey"></td>
388 <td></td>
389 <td class="paramtype">uint&#160;</td>
390 <td class="paramname"><em>src_stride_x</em>, </td>
391 </tr>
392 <tr>
393 <td class="paramkey"></td>
394 <td></td>
395 <td class="paramtype">uint&#160;</td>
396 <td class="paramname"><em>src_step_x</em>, </td>
397 </tr>
398 <tr>
399 <td class="paramkey"></td>
400 <td></td>
401 <td class="paramtype">uint&#160;</td>
402 <td class="paramname"><em>src_stride_y</em>, </td>
403 </tr>
404 <tr>
405 <td class="paramkey"></td>
406 <td></td>
407 <td class="paramtype">uint&#160;</td>
408 <td class="paramname"><em>src_step_y</em>, </td>
409 </tr>
410 <tr>
411 <td class="paramkey"></td>
412 <td></td>
413 <td class="paramtype">uint&#160;</td>
414 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
415 </tr>
416 <tr>
417 <td class="paramkey"></td>
418 <td></td>
419 <td class="paramtype">__global uchar *&#160;</td>
420 <td class="paramname"><em>dst_ptr</em>, </td>
421 </tr>
422 <tr>
423 <td class="paramkey"></td>
424 <td></td>
425 <td class="paramtype">uint&#160;</td>
426 <td class="paramname"><em>dst_stride_x</em>, </td>
427 </tr>
428 <tr>
429 <td class="paramkey"></td>
430 <td></td>
431 <td class="paramtype">uint&#160;</td>
432 <td class="paramname"><em>dst_step_x</em>, </td>
433 </tr>
434 <tr>
435 <td class="paramkey"></td>
436 <td></td>
437 <td class="paramtype">uint&#160;</td>
438 <td class="paramname"><em>dst_stride_y</em>, </td>
439 </tr>
440 <tr>
441 <td class="paramkey"></td>
442 <td></td>
443 <td class="paramtype">uint&#160;</td>
444 <td class="paramname"><em>dst_step_y</em>, </td>
445 </tr>
446 <tr>
447 <td class="paramkey"></td>
448 <td></td>
449 <td class="paramtype">uint&#160;</td>
450 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
451 </tr>
452 <tr>
453 <td></td>
454 <td>)</td>
455 <td></td><td></td>
456 </tr>
457 </table>
458</div><div class="memdoc">
459
460<p>This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values. </p>
461<dl class="params"><dt>Parameters</dt><dd>
462 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]463 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U8/S8/QS8 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]464 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
465 <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>
466 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
467 <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>
468 <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 matrix </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]469 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]470 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
471 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
472 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
473 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
474 <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 matrix </td></tr>
475 </table>
476 </dd>
477</dl>
478
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]479<p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00223">223</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]480
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]481<p>References <a class="el" href="fixed__point_8h_source.xhtml#l00178">ADD_SAT_OP_EXPAND</a>, <a class="el" href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00117">arm_compute::test::validation::c</a>, <a class="el" href="helpers_8h_source.xhtml#l00102">CONVERT_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00096">CONVERT_TO_VECTOR_STRUCT</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">arm_compute::test::validation::dst</a>, <a class="el" href="helpers_8h_source.xhtml#l00066">IMAGE_DECLARATION</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00276">mla_sat_qs16x8()</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00275">mla_sat_qs8x16()</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00296">mlal_sat_qs16x8()</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00295">mlal_sat_qs8x8()</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00254">mul_sat_qs16x8()</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00249">mul_sat_qs8x16()</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#l00074">TENSOR3D_DECLARATION</a>, <a class="el" href="fixed__point_8h_source.xhtml#l00093">VEC_DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00060">VECTOR_DECLARATION</a>, <a class="el" href="softmax__layer_8cl_source.xhtml#l00079">VECTOR_SIZE</a>, <a class="el" href="helpers_8h_source.xhtml#l00040">VLOAD</a>, and <a class="el" href="helpers_8h_source.xhtml#l00043">VSTORE</a>.</p>
482<div class="fragment"><div class="line"><a name="l00225"></a><span class="lineno"> 225</span>&#160;{</div><div class="line"><a name="l00226"></a><span class="lineno"> 226</span>&#160; <span class="comment">// Compute source and destination addresses</span></div><div class="line"><a name="l00227"></a><span class="lineno"> 227</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00228"></a><span class="lineno"> 228</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#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00229"></a><span class="lineno"> 229</span>&#160;</div><div class="line"><a name="l00230"></a><span class="lineno"> 230</span>&#160; <span class="comment">// Load values from Matrix A</span></div><div class="line"><a name="l00231"></a><span class="lineno"> 231</span>&#160; uchar16 a0 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 0)));</div><div class="line"><a name="l00232"></a><span class="lineno"> 232</span>&#160; uchar16 a1 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 1)));</div><div class="line"><a name="l00233"></a><span class="lineno"> 233</span>&#160; uchar16 a2 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 2)));</div><div class="line"><a name="l00234"></a><span class="lineno"> 234</span>&#160; uchar16 a3 = vload16(0, (__global uchar *)(<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&amp;src, 0, 3)));</div><div class="line"><a name="l00235"></a><span class="lineno"> 235</span>&#160;</div><div class="line"><a name="l00236"></a><span class="lineno"> 236</span>&#160; uchar16 val0 = (uchar16)((uchar4)(a0.s0, a1.s0, a2.s0, a3.s0), (uchar4)(a0.s1, a1.s1, a2.s1, a3.s1),</div><div class="line"><a name="l00237"></a><span class="lineno"> 237</span>&#160; (uchar4)(a0.s2, a1.s2, a2.s2, a3.s2), (uchar4)(a0.s3, a1.s3, a2.s3, a3.s3));</div><div class="line"><a name="l00238"></a><span class="lineno"> 238</span>&#160; vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 0);</div><div class="line"><a name="l00239"></a><span class="lineno"> 239</span>&#160;</div><div class="line"><a name="l00240"></a><span class="lineno"> 240</span>&#160; val0 = (uchar16)((uchar4)(a0.s4, a1.s4, a2.s4, a3.s4), (uchar4)(a0.s5, a1.s5, a2.s5, a3.s5),</div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>&#160; (uchar4)(a0.s6, a1.s6, a2.s6, a3.s6), (uchar4)(a0.s7, a1.s7, a2.s7, a3.s7));</div><div class="line"><a name="l00242"></a><span class="lineno"> 242</span>&#160; vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 16);</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span>&#160;</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span>&#160; val0 = (uchar16)((uchar4)(a0.s8, a1.s8, a2.s8, a3.s8), (uchar4)(a0.s9, a1.s9, a2.s9, a3.s9),</div><div class="line"><a name="l00245"></a><span class="lineno"> 245</span>&#160; (uchar4)(a0.sA, a1.sA, a2.sA, a3.sA), (uchar4)(a0.sB, a1.sB, a2.sB, a3.sB));</div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span>&#160; vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 32);</div><div class="line"><a name="l00247"></a><span class="lineno"> 247</span>&#160;</div><div class="line"><a name="l00248"></a><span class="lineno"> 248</span>&#160; val0 = (uchar16)((uchar4)(a0.sC, a1.sC, a2.sC, a3.sC), (uchar4)(a0.sD, a1.sD, a2.sD, a3.sD),</div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>&#160; (uchar4)(a0.sE, a1.sE, a2.sE, a3.sE), (uchar4)(a0.sF, a1.sF, a2.sF, a3.sF));</div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>&#160; vstore16(val0, 0, ((__global uchar *)dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) + 48);</div><div class="line"><a name="l00251"></a><span class="lineno"> 251</span>&#160;}</div><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>
483<div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00102">helpers.h:102</a></div></div>
484<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">Tensor dst</div><div class="ttdef"><b>Definition:</b> <a href="validation_2_n_e_o_n_2_g_e_m_m_8cpp_source.xhtml#l00118">GEMM.cpp:118</a></div></div>
485<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>
486<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]487</div><!-- fragment -->
488</div>
489</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]490<a class="anchor" id="a97a299ffb620aa98fe5e39bb8ab0833f"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]491<div class="memitem">
492<div class="memproto">
493 <table class="memname">
494 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]495 <td class="memname">__kernel void gemm_transpose1x16 </td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]496 <td>(</td>
497 <td class="paramtype">__global uchar *&#160;</td>
498 <td class="paramname"><em>src_ptr</em>, </td>
499 </tr>
500 <tr>
501 <td class="paramkey"></td>
502 <td></td>
503 <td class="paramtype">uint&#160;</td>
504 <td class="paramname"><em>src_stride_x</em>, </td>
505 </tr>
506 <tr>
507 <td class="paramkey"></td>
508 <td></td>
509 <td class="paramtype">uint&#160;</td>
510 <td class="paramname"><em>src_step_x</em>, </td>
511 </tr>
512 <tr>
513 <td class="paramkey"></td>
514 <td></td>
515 <td class="paramtype">uint&#160;</td>
516 <td class="paramname"><em>src_stride_y</em>, </td>
517 </tr>
518 <tr>
519 <td class="paramkey"></td>
520 <td></td>
521 <td class="paramtype">uint&#160;</td>
522 <td class="paramname"><em>src_step_y</em>, </td>
523 </tr>
524 <tr>
525 <td class="paramkey"></td>
526 <td></td>
527 <td class="paramtype">uint&#160;</td>
528 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
529 </tr>
530 <tr>
531 <td class="paramkey"></td>
532 <td></td>
533 <td class="paramtype">__global uchar *&#160;</td>
534 <td class="paramname"><em>dst_ptr</em>, </td>
535 </tr>
536 <tr>
537 <td class="paramkey"></td>
538 <td></td>
539 <td class="paramtype">uint&#160;</td>
540 <td class="paramname"><em>dst_stride_x</em>, </td>
541 </tr>
542 <tr>
543 <td class="paramkey"></td>
544 <td></td>
545 <td class="paramtype">uint&#160;</td>
546 <td class="paramname"><em>dst_step_x</em>, </td>
547 </tr>
548 <tr>
549 <td class="paramkey"></td>
550 <td></td>
551 <td class="paramtype">uint&#160;</td>
552 <td class="paramname"><em>dst_stride_y</em>, </td>
553 </tr>
554 <tr>
555 <td class="paramkey"></td>
556 <td></td>
557 <td class="paramtype">uint&#160;</td>
558 <td class="paramname"><em>dst_step_y</em>, </td>
559 </tr>
560 <tr>
561 <td class="paramkey"></td>
562 <td></td>
563 <td class="paramtype">uint&#160;</td>
564 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
565 </tr>
566 <tr>
567 <td></td>
568 <td>)</td>
569 <td></td><td></td>
570 </tr>
571 </table>
572</div><div class="memdoc">
573
574<p>This OpenCL kernel computes the "vector" 1x16 transposition of input matrix. </p>
575<dl class="params"><dt>Parameters</dt><dd>
576 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]577 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U8/S8/QS8 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]578 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
579 <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>
580 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
581 <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>
582 <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 matrix </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]583 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]584 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
585 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
586 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
587 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
588 <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 matrix </td></tr>
589 </table>
590 </dd>
591</dl>
592
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]593<p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00109">109</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]594
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]595<p>References <a class="el" href="helpers_8h_source.xhtml#l00102">CONVERT_TO_IMAGE_STRUCT</a>, and <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>.</p>
596<div class="fragment"><div class="line"><a name="l00111"></a><span class="lineno"> 111</span>&#160;{</div><div class="line"><a name="l00112"></a><span class="lineno"> 112</span>&#160; uint x = get_global_id(0);</div><div class="line"><a name="l00113"></a><span class="lineno"> 113</span>&#160; uint y = get_global_id(1);</div><div class="line"><a name="l00114"></a><span class="lineno"> 114</span>&#160;</div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span>&#160; <span class="comment">// Compute address for Matrix B - source</span></div><div class="line"><a name="l00116"></a><span class="lineno"> 116</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00117"></a><span class="lineno"> 117</span>&#160;</div><div class="line"><a name="l00118"></a><span class="lineno"> 118</span>&#160; <span class="comment">// Compute address for Matrix B transposed - destination. X and Y are swapped</span></div><div class="line"><a name="l00119"></a><span class="lineno"> 119</span>&#160; uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));</div><div class="line"><a name="l00120"></a><span class="lineno"> 120</span>&#160;</div><div class="line"><a name="l00121"></a><span class="lineno"> 121</span>&#160; uchar16 b0 = vload16(0, (__global uchar *)src.ptr);</div><div class="line"><a name="l00122"></a><span class="lineno"> 122</span>&#160;</div><div class="line"><a name="l00123"></a><span class="lineno"> 123</span>&#160; vstore16(b0, 0, (__global uchar *)(dst_ptr + dst_addr_in_bytes));</div><div class="line"><a name="l00124"></a><span class="lineno"> 124</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00102">helpers.h:102</a></div></div>
597<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]598</div><!-- fragment -->
599</div>
600</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]601<a class="anchor" id="ac5eba21b9ac015e1454044aeba93546c"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]602<div class="memitem">
603<div class="memproto">
604 <table class="memname">
605 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]606 <td class="memname">__kernel void gemm_transpose1x4 </td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]607 <td>(</td>
608 <td class="paramtype">__global uchar *&#160;</td>
609 <td class="paramname"><em>src_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>src_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>src_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>src_stride_y</em>, </td>
628 </tr>
629 <tr>
630 <td class="paramkey"></td>
631 <td></td>
632 <td class="paramtype">uint&#160;</td>
633 <td class="paramname"><em>src_step_y</em>, </td>
634 </tr>
635 <tr>
636 <td class="paramkey"></td>
637 <td></td>
638 <td class="paramtype">uint&#160;</td>
639 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
640 </tr>
641 <tr>
642 <td class="paramkey"></td>
643 <td></td>
644 <td class="paramtype">__global uchar *&#160;</td>
645 <td class="paramname"><em>dst_ptr</em>, </td>
646 </tr>
647 <tr>
648 <td class="paramkey"></td>
649 <td></td>
650 <td class="paramtype">uint&#160;</td>
651 <td class="paramname"><em>dst_stride_x</em>, </td>
652 </tr>
653 <tr>
654 <td class="paramkey"></td>
655 <td></td>
656 <td class="paramtype">uint&#160;</td>
657 <td class="paramname"><em>dst_step_x</em>, </td>
658 </tr>
659 <tr>
660 <td class="paramkey"></td>
661 <td></td>
662 <td class="paramtype">uint&#160;</td>
663 <td class="paramname"><em>dst_stride_y</em>, </td>
664 </tr>
665 <tr>
666 <td class="paramkey"></td>
667 <td></td>
668 <td class="paramtype">uint&#160;</td>
669 <td class="paramname"><em>dst_step_y</em>, </td>
670 </tr>
671 <tr>
672 <td class="paramkey"></td>
673 <td></td>
674 <td class="paramtype">uint&#160;</td>
675 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
676 </tr>
677 <tr>
678 <td></td>
679 <td>)</td>
680 <td></td><td></td>
681 </tr>
682 </table>
683</div><div class="memdoc">
684
685<p>This OpenCL kernel computes the "vector" 1x4 transposition of input matrix. </p>
686<dl class="params"><dt>Parameters</dt><dd>
687 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]688 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U32/S32/F32 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]689 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
690 <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>
691 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
692 <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>
693 <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 matrix </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]694 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]695 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
696 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
697 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
698 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
699 <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 matrix </td></tr>
700 </table>
701 </dd>
702</dl>
703
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]704<p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00045">45</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]705
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]706<p>References <a class="el" href="helpers_8h_source.xhtml#l00102">CONVERT_TO_IMAGE_STRUCT</a>, and <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>.</p>
707<div class="fragment"><div class="line"><a name="l00047"></a><span class="lineno"> 47</span>&#160;{</div><div class="line"><a name="l00048"></a><span class="lineno"> 48</span>&#160; uint x = get_global_id(0);</div><div class="line"><a name="l00049"></a><span class="lineno"> 49</span>&#160; uint y = get_global_id(1);</div><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; <span class="comment">// Compute address for Matrix B - source</span></div><div class="line"><a name="l00052"></a><span class="lineno"> 52</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</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; <span class="comment">// Compute address for Matrix B transposed - destination. X and Y are swapped</span></div><div class="line"><a name="l00055"></a><span class="lineno"> 55</span>&#160; uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));</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; uint4 b0 = vload4(0, (__global uint *)src.ptr);</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>&#160;</div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>&#160; vstore4(b0, 0, (__global uint *)(dst_ptr + dst_addr_in_bytes));</div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00102">helpers.h:102</a></div></div>
708<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]709</div><!-- fragment -->
710</div>
711</div>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]712<a class="anchor" id="a75c35ac120f123e9cce56c43ba3f48d2"></a>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]713<div class="memitem">
714<div class="memproto">
715 <table class="memname">
716 <tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]717 <td class="memname">__kernel void gemm_transpose1x8 </td>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]718 <td>(</td>
719 <td class="paramtype">__global uchar *&#160;</td>
720 <td class="paramname"><em>src_ptr</em>, </td>
721 </tr>
722 <tr>
723 <td class="paramkey"></td>
724 <td></td>
725 <td class="paramtype">uint&#160;</td>
726 <td class="paramname"><em>src_stride_x</em>, </td>
727 </tr>
728 <tr>
729 <td class="paramkey"></td>
730 <td></td>
731 <td class="paramtype">uint&#160;</td>
732 <td class="paramname"><em>src_step_x</em>, </td>
733 </tr>
734 <tr>
735 <td class="paramkey"></td>
736 <td></td>
737 <td class="paramtype">uint&#160;</td>
738 <td class="paramname"><em>src_stride_y</em>, </td>
739 </tr>
740 <tr>
741 <td class="paramkey"></td>
742 <td></td>
743 <td class="paramtype">uint&#160;</td>
744 <td class="paramname"><em>src_step_y</em>, </td>
745 </tr>
746 <tr>
747 <td class="paramkey"></td>
748 <td></td>
749 <td class="paramtype">uint&#160;</td>
750 <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
751 </tr>
752 <tr>
753 <td class="paramkey"></td>
754 <td></td>
755 <td class="paramtype">__global uchar *&#160;</td>
756 <td class="paramname"><em>dst_ptr</em>, </td>
757 </tr>
758 <tr>
759 <td class="paramkey"></td>
760 <td></td>
761 <td class="paramtype">uint&#160;</td>
762 <td class="paramname"><em>dst_stride_x</em>, </td>
763 </tr>
764 <tr>
765 <td class="paramkey"></td>
766 <td></td>
767 <td class="paramtype">uint&#160;</td>
768 <td class="paramname"><em>dst_step_x</em>, </td>
769 </tr>
770 <tr>
771 <td class="paramkey"></td>
772 <td></td>
773 <td class="paramtype">uint&#160;</td>
774 <td class="paramname"><em>dst_stride_y</em>, </td>
775 </tr>
776 <tr>
777 <td class="paramkey"></td>
778 <td></td>
779 <td class="paramtype">uint&#160;</td>
780 <td class="paramname"><em>dst_step_y</em>, </td>
781 </tr>
782 <tr>
783 <td class="paramkey"></td>
784 <td></td>
785 <td class="paramtype">uint&#160;</td>
786 <td class="paramname"><em>dst_offset_first_element_in_bytes</em>&#160;</td>
787 </tr>
788 <tr>
789 <td></td>
790 <td>)</td>
791 <td></td><td></td>
792 </tr>
793 </table>
794</div><div class="memdoc">
795
796<p>This OpenCL kernel computes the "vector" 1x8 transposition of input matrix. </p>
797<dl class="params"><dt>Parameters</dt><dd>
798 <table class="params">
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]799 <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source matrix. Supported data types: U16/S16/QS16/F16 </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]800 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source matrix in X dimension (in bytes) </td></tr>
801 <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>
802 <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source matrix in Y dimension (in bytes) </td></tr>
803 <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>
804 <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 matrix </td></tr>
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]805 <tr><td class="paramdir">[out]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination matrix Supported data types: same as <code>src_ptr</code> </td></tr>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]806 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_x</td><td>Stride of the destination matrix in X dimension (in bytes) </td></tr>
807 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_x</td><td>dst_gx_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
808 <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_y</td><td>Stride of the destination matrix in Y dimension (in bytes) </td></tr>
809 <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_y</td><td>dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
810 <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 matrix </td></tr>
811 </table>
812 </dd>
813</dl>
814
Kaizen8938bd32017-09-28 14:38:23 +0100[diff] [blame]815<p>Definition at line <a class="el" href="gemm_8cl_source.xhtml#l00077">77</a> of file <a class="el" href="gemm_8cl_source.xhtml">gemm.cl</a>.</p>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]816
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]817<p>References <a class="el" href="helpers_8h_source.xhtml#l00102">CONVERT_TO_IMAGE_STRUCT</a>, and <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>.</p>
818<div class="fragment"><div class="line"><a name="l00079"></a><span class="lineno"> 79</span>&#160;{</div><div class="line"><a name="l00080"></a><span class="lineno"> 80</span>&#160; uint x = get_global_id(0);</div><div class="line"><a name="l00081"></a><span class="lineno"> 81</span>&#160; uint y = get_global_id(1);</div><div class="line"><a name="l00082"></a><span class="lineno"> 82</span>&#160;</div><div class="line"><a name="l00083"></a><span class="lineno"> 83</span>&#160; <span class="comment">// Compute address for Matrix B - source</span></div><div class="line"><a name="l00084"></a><span class="lineno"> 84</span>&#160; <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00085"></a><span class="lineno"> 85</span>&#160;</div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span>&#160; <span class="comment">// Compute address for Matrix B transposed - destination. X and Y are swapped</span></div><div class="line"><a name="l00087"></a><span class="lineno"> 87</span>&#160; uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));</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; ushort8 b0 = vload8(0, (__global ushort *)src.ptr);</div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>&#160;</div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span>&#160; vstore8(b0, 0, (__global ushort *)(dst_ptr + dst_addr_in_bytes));</div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span>&#160;}</div><div class="ttc" id="helpers_8h_xhtml_aebe814363556c244be043b13e7969197"><div class="ttname"><a href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00102">helpers.h:102</a></div></div>
819<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>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]820</div><!-- fragment -->
821</div>
822</div>
823</div><!-- contents -->
824</div><!-- doc-content -->
825<!-- start footer part -->
826<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
827 <ul>
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]828 <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="gemm_8cl.xhtml">gemm.cl</a></li>
829 <li class="footer">Generated on Thu Dec 14 2017 23:48:34 for Compute Library by
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]830 <a href="http://www.doxygen.org/index.html">
Anthony Barbier8140e1e2017-12-14 23:48:46 +0000[diff] [blame^]831 <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
Anthony Barbier871448e2017-03-24 14:54:29 +0000[diff] [blame]832 </ul>
833</div>
834</body>
835</html>