arm_compute v18.02
Change-Id: I7207aa488e5470f235f39b6c188b4678dc38d1a6
diff --git a/documentation/depthwise__convolution_8cl.xhtml b/documentation/depthwise__convolution_8cl.xhtml
index 2bd7be7..0a7875d 100644
--- a/documentation/depthwise__convolution_8cl.xhtml
+++ b/documentation/depthwise__convolution_8cl.xhtml
@@ -40,7 +40,7 @@
<tr style="height: 56px;">
<td style="padding-left: 0.5em;">
<div id="projectname">Compute Library
-  <span id="projectnumber">18.01</span>
+  <span id="projectnumber">18.02</span>
</div>
</td>
</tr>
@@ -113,6 +113,9 @@
</div>
<div class="header">
+ <div class="summary">
+<a href="#define-members">Macros</a> |
+<a href="#func-members">Functions</a> </div>
<div class="headertitle">
<div class="title">depthwise_convolution.cl File Reference</div> </div>
</div><!--header-->
@@ -120,13 +123,578 @@
<div class="textblock"><code>#include "<a class="el" href="helpers_8h_source.xhtml">helpers.h</a>"</code><br />
</div>
<p><a href="depthwise__convolution_8cl_source.xhtml">Go to the source code of this file.</a></p>
+<table class="memberdecls">
+<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="define-members"></a>
+Macros</h2></td></tr>
+<tr class="memitem:aa5ab4cf85fe56de3821727cfcef327d9"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(acc, src0, weights_row0)</td></tr>
+<tr class="separator:aa5ab4cf85fe56de3821727cfcef327d9"><td class="memSeparator" colspan="2"> </td></tr>
+<tr class="memitem:a31cbdd6b28ee6331839432f69d87fdcb"><td class="memItemLeft" align="right" valign="top">#define </td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(acc, src0, src1, weights_row0)</td></tr>
+<tr class="separator:a31cbdd6b28ee6331839432f69d87fdcb"><td class="memSeparator" colspan="2"> </td></tr>
+</table><table class="memberdecls">
+<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
+Functions</h2></td></tr>
+<tr class="memitem:a2fbbf5b73260b66b310d4d857db5289e"><td class="memItemLeft" align="right" valign="top">__kernel void </td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a2fbbf5b73260b66b310d4d857db5289e">depthwise_convolution_3x3_stridex1_stridey1_bifrost</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_stride_z, uint src_step_z, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes, __global uchar *weights_ptr, uint weights_stride_x, uint weights_step_x, uint weights_stride_y, uint weights_step_y, uint weights_stride_z, uint weights_step_z, uint weights_offset_first_element_in_bytes, __global uchar *biases_ptr, uint biases_stride_x, uint biases_step_x, uint biases_offset_first_element_in_bytes)</td></tr>
+<tr class="memdesc:a2fbbf5b73260b66b310d4d857db5289e"><td class="mdescLeft"> </td><td class="mdescRight">This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 1. <a href="#a2fbbf5b73260b66b310d4d857db5289e">More...</a><br /></td></tr>
+<tr class="separator:a2fbbf5b73260b66b310d4d857db5289e"><td class="memSeparator" colspan="2"> </td></tr>
+<tr class="memitem:a328ab548c536089e8238fb33de617bec"><td class="memItemLeft" align="right" valign="top">__kernel void </td><td class="memItemRight" valign="bottom"><a class="el" href="depthwise__convolution_8cl.xhtml#a328ab548c536089e8238fb33de617bec">depthwise_convolution_3x3_stridex2_stridey2_bifrost</a> (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_stride_z, uint src_step_z, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes, __global uchar *weights_ptr, uint weights_stride_x, uint weights_step_x, uint weights_stride_y, uint weights_step_y, uint weights_stride_z, uint weights_step_z, uint weights_offset_first_element_in_bytes, __global uchar *biases_ptr, uint biases_stride_x, uint biases_step_x, uint biases_offset_first_element_in_bytes)</td></tr>
+<tr class="memdesc:a328ab548c536089e8238fb33de617bec"><td class="mdescLeft"> </td><td class="mdescRight">This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 2. <a href="#a328ab548c536089e8238fb33de617bec">More...</a><br /></td></tr>
+<tr class="separator:a328ab548c536089e8238fb33de617bec"><td class="memSeparator" colspan="2"> </td></tr>
+</table>
+<h2 class="groupheader">Macro Definition Documentation</h2>
+<a class="anchor" id="aa5ab4cf85fe56de3821727cfcef327d9"></a>
+<div class="memitem">
+<div class="memproto">
+ <table class="memname">
+ <tr>
+ <td class="memname">#define CONVOLUTION1x3_BIFROST2X1_STRIDE1</td>
+ <td>(</td>
+ <td class="paramtype"> </td>
+ <td class="paramname">acc, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype"> </td>
+ <td class="paramname">src0, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype"> </td>
+ <td class="paramname">weights_row0 </td>
+ </tr>
+ <tr>
+ <td></td>
+ <td>)</td>
+ <td></td><td></td>
+ </tr>
+ </table>
+</div><div class="memdoc">
+<b>Value:</b><div class="fragment"><div class="line">({ \</div><div class="line"> acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0); \</div><div class="line"> acc.s1 = fma(src0.s1, weights_row0.s0, acc.s1); \</div><div class="line"> acc.s1 = fma(src0.s2, weights_row0.s1, acc.s1); \</div><div class="line"> acc.s1 = fma(src0.s3, weights_row0.s2, acc.s1); \</div><div class="line"> })</div></div><!-- fragment -->
+<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00211">211</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
+
+<p>Referenced by <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00263">depthwise_convolution_3x3_stridex1_stridey1_bifrost()</a>.</p>
+
+</div>
+</div>
+<a class="anchor" id="a31cbdd6b28ee6331839432f69d87fdcb"></a>
+<div class="memitem">
+<div class="memproto">
+ <table class="memname">
+ <tr>
+ <td class="memname">#define CONVOLUTION1x3_BIFROST2X1_STRIDE2</td>
+ <td>(</td>
+ <td class="paramtype"> </td>
+ <td class="paramname">acc, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype"> </td>
+ <td class="paramname">src0, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype"> </td>
+ <td class="paramname">src1, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype"> </td>
+ <td class="paramname">weights_row0 </td>
+ </tr>
+ <tr>
+ <td></td>
+ <td>)</td>
+ <td></td><td></td>
+ </tr>
+ </table>
+</div><div class="memdoc">
+<b>Value:</b><div class="fragment"><div class="line">({ \</div><div class="line"> acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0); \</div><div class="line"> acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0); \</div><div class="line"> acc.s1 = fma(src0.s2, weights_row0.s0, acc.s1); \</div><div class="line"> acc.s1 = fma(src0.s3, weights_row0.s1, acc.s1); \</div><div class="line"> acc.s1 = fma(src1.s0, weights_row0.s2, acc.s1); \</div><div class="line"> })</div></div><!-- fragment -->
+<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00221">221</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
+
+<p>Referenced by <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00360">depthwise_convolution_3x3_stridex2_stridey2_bifrost()</a>.</p>
+
+</div>
+</div>
+<h2 class="groupheader">Function Documentation</h2>
+<a class="anchor" id="a2fbbf5b73260b66b310d4d857db5289e"></a>
+<div class="memitem">
+<div class="memproto">
+ <table class="memname">
+ <tr>
+ <td class="memname">__kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost </td>
+ <td>(</td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>src_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>dst_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>weights_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_offset_first_element_in_bytes</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>biases_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>biases_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>biases_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>biases_offset_first_element_in_bytes</em> </td>
+ </tr>
+ <tr>
+ <td></td>
+ <td>)</td>
+ <td></td><td></td>
+ </tr>
+ </table>
+</div><div class="memdoc">
+
+<p>This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 1. </p>
+<dl class="params"><dt>Parameters</dt><dd>
+ <table class="params">
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source image. Supported data types: F32 </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_offset_first_element_in_bytes</td><td>The offset of the first element in the source image </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F32 </td></tr>
+ <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>
+ <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>
+ <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>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_ptr</td><td>Pointer to the weights tensor. Supported data types: F32 </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_x</td><td>Stride of the weights tensor in X dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_x</td><td>weights_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_y</td><td>Stride of the weights tensor in Y dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_y</td><td>weights_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_z</td><td>Stride of the weights tensor in Z dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_z</td><td>weights_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_offset_first_element_in_bytes</td><td>The offset of the first element in the biases vector </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>(Optional) Pointer to the biases vector. Supported data types: F32 </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>(Optional) Stride of the biases vector in X dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>(Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>(Optional) The offset of the first element in the biases vector </td></tr>
+ </table>
+ </dd>
+</dl>
+
+<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00263">263</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
+
+<p>References <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00117">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00099">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>, <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00211">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00151">Tensor3D::ptr</a>, and <a class="el" href="helpers_8h_source.xhtml#l00290">vector_offset()</a>.</p>
+<div class="fragment"><div class="line"><a name="l00272"></a><span class="lineno"> 272</span> {</div><div class="line"><a name="l00273"></a><span class="lineno"> 273</span>  <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00274"></a><span class="lineno"> 274</span>  <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00275"></a><span class="lineno"> 275</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> weights = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(weights);</div><div class="line"><a name="l00276"></a><span class="lineno"> 276</span> </div><div class="line"><a name="l00277"></a><span class="lineno"> 277</span>  float2 pixels0 = 0.0f;</div><div class="line"><a name="l00278"></a><span class="lineno"> 278</span>  float2 pixels1 = 0.0f;</div><div class="line"><a name="l00279"></a><span class="lineno"> 279</span>  float2 pixels2 = 0.0f;</div><div class="line"><a name="l00280"></a><span class="lineno"> 280</span>  float2 pixels3 = 0.0f;</div><div class="line"><a name="l00281"></a><span class="lineno"> 281</span> </div><div class="line"><a name="l00282"></a><span class="lineno"> 282</span>  __global uchar *weights_addr = (__global uchar *)weights.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>;</div><div class="line"><a name="l00283"></a><span class="lineno"> 283</span>  __global uchar *src_addr = (__global uchar *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, 0, 0);</div><div class="line"><a name="l00284"></a><span class="lineno"> 284</span> </div><div class="line"><a name="l00285"></a><span class="lineno"> 285</span>  <span class="comment">// Load the weights</span></div><div class="line"><a name="l00286"></a><span class="lineno"> 286</span>  float3 weights_row0 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 0 * weights_stride_y));</div><div class="line"><a name="l00287"></a><span class="lineno"> 287</span>  float3 weights_row1 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 1 * weights_stride_y));</div><div class="line"><a name="l00288"></a><span class="lineno"> 288</span>  float3 weights_row2 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 2 * weights_stride_y));</div><div class="line"><a name="l00289"></a><span class="lineno"> 289</span> </div><div class="line"><a name="l00290"></a><span class="lineno"> 290</span>  <span class="comment">// Note: Since each work-item computes 4x2 elements, we need to load 4 rows from the input tensor</span></div><div class="line"><a name="l00291"></a><span class="lineno"> 291</span>  float4 src00 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00292"></a><span class="lineno"> 292</span>  float4 src10 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00293"></a><span class="lineno"> 293</span>  float4 src20 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00294"></a><span class="lineno"> 294</span>  float4 src30 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00295"></a><span class="lineno"> 295</span>  float4 src40 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00296"></a><span class="lineno"> 296</span>  float4 src50 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 5 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00297"></a><span class="lineno"> 297</span> </div><div class="line"><a name="l00298"></a><span class="lineno"> 298</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src00, weights_row0);</div><div class="line"><a name="l00299"></a><span class="lineno"> 299</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src10, weights_row1);</div><div class="line"><a name="l00300"></a><span class="lineno"> 300</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels0, src20, weights_row2);</div><div class="line"><a name="l00301"></a><span class="lineno"> 301</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src10, weights_row0);</div><div class="line"><a name="l00302"></a><span class="lineno"> 302</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src20, weights_row1);</div><div class="line"><a name="l00303"></a><span class="lineno"> 303</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels1, src30, weights_row2);</div><div class="line"><a name="l00304"></a><span class="lineno"> 304</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src20, weights_row0);</div><div class="line"><a name="l00305"></a><span class="lineno"> 305</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src30, weights_row1);</div><div class="line"><a name="l00306"></a><span class="lineno"> 306</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels2, src40, weights_row2);</div><div class="line"><a name="l00307"></a><span class="lineno"> 307</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src30, weights_row0);</div><div class="line"><a name="l00308"></a><span class="lineno"> 308</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src40, weights_row1);</div><div class="line"><a name="l00309"></a><span class="lineno"> 309</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a>(pixels3, src50, weights_row2);</div><div class="line"><a name="l00310"></a><span class="lineno"> 310</span> </div><div class="line"><a name="l00311"></a><span class="lineno"> 311</span> <span class="preprocessor">#ifdef HAS_BIAS</span></div><div class="line"><a name="l00312"></a><span class="lineno"> 312</span>  <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>(biases);</div><div class="line"><a name="l00313"></a><span class="lineno"> 313</span> </div><div class="line"><a name="l00314"></a><span class="lineno"> 314</span>  <span class="keywordtype">float</span> bias = *((__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a>(&biases, get_global_id(2))));</div><div class="line"><a name="l00315"></a><span class="lineno"> 315</span> </div><div class="line"><a name="l00316"></a><span class="lineno"> 316</span>  pixels0 += (float2)bias;</div><div class="line"><a name="l00317"></a><span class="lineno"> 317</span>  pixels1 += (float2)bias;</div><div class="line"><a name="l00318"></a><span class="lineno"> 318</span>  pixels2 += (float2)bias;</div><div class="line"><a name="l00319"></a><span class="lineno"> 319</span>  pixels3 += (float2)bias;</div><div class="line"><a name="l00320"></a><span class="lineno"> 320</span> <span class="preprocessor">#endif </span><span class="comment">/* defined(HAS_BIAS) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00321"></a><span class="lineno"> 321</span> </div><div class="line"><a name="l00322"></a><span class="lineno"> 322</span>  vstore2(pixels0, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 0 * dst_stride_y));</div><div class="line"><a name="l00323"></a><span class="lineno"> 323</span>  vstore2(pixels1, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 1 * dst_stride_y));</div><div class="line"><a name="l00324"></a><span class="lineno"> 324</span>  vstore2(pixels2, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 2 * dst_stride_y));</div><div class="line"><a name="l00325"></a><span class="lineno"> 325</span>  vstore2(pixels3, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 3 * dst_stride_y));</div><div class="line"><a name="l00326"></a><span class="lineno"> 326</span> }</div><div class="ttc" id="struct_vector_xhtml"><div class="ttname"><a href="struct_vector.xhtml">Vector</a></div><div class="ttdoc">Structure to hold Vector information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00132">helpers.h:132</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00117">helpers.h:117</a></div></div>
+<div class="ttc" id="struct_tensor3_d_xhtml"><div class="ttname"><a href="struct_tensor3_d.xhtml">Tensor3D</a></div><div class="ttdoc">Structure to hold 3D tensor information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00149">helpers.h:149</a></div></div>
+<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>
+<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>
+<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>
+<div class="ttc" id="depthwise__convolution_8cl_xhtml_aa5ab4cf85fe56de3821727cfcef327d9"><div class="ttname"><a href="depthwise__convolution_8cl.xhtml#aa5ab4cf85fe56de3821727cfcef327d9">CONVOLUTION1x3_BIFROST2X1_STRIDE1</a></div><div class="ttdeci">#define CONVOLUTION1x3_BIFROST2X1_STRIDE1(acc, src0, weights_row0)</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution_8cl_source.xhtml#l00211">depthwise_convolution.cl:211</a></div></div>
+<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>
+<div class="ttc" id="helpers_8h_xhtml_a64d779f80eeb923e0ab2313433f7b40b"><div class="ttname"><a href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a7e4940407322d6f0ccb8b6b86b856019"><div class="ttname"><a href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a></div><div class="ttdeci">__global const uchar * vector_offset(const Vector *vec, int x)</div><div class="ttdoc">Get the pointer position of a Vector. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00290">helpers.h:290</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
+<div class="ttc" id="struct_tensor3_d_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">Tensor3D::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00151">helpers.h:151</a></div></div>
+<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 & src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
+</div><!-- fragment -->
+</div>
+</div>
+<a class="anchor" id="a328ab548c536089e8238fb33de617bec"></a>
+<div class="memitem">
+<div class="memproto">
+ <table class="memname">
+ <tr>
+ <td class="memname">__kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost </td>
+ <td>(</td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>src_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>src_offset_first_element_in_bytes</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>dst_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>dst_offset_first_element_in_bytes</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>weights_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_stride_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_step_y</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_stride_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_step_z</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>weights_offset_first_element_in_bytes</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">__global uchar * </td>
+ <td class="paramname"><em>biases_ptr</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>biases_stride_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>biases_step_x</em>, </td>
+ </tr>
+ <tr>
+ <td class="paramkey"></td>
+ <td></td>
+ <td class="paramtype">uint </td>
+ <td class="paramname"><em>biases_offset_first_element_in_bytes</em> </td>
+ </tr>
+ <tr>
+ <td></td>
+ <td>)</td>
+ <td></td><td></td>
+ </tr>
+ </table>
+</div><div class="memdoc">
+
+<p>This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both stride_x and stride_y are equal to 2. </p>
+<dl class="params"><dt>Parameters</dt><dd>
+ <table class="params">
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_ptr</td><td>Pointer to the source image. Supported data types: F32 </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_x</td><td>Stride of the source image in X dimension (in bytes) </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_stride_y</td><td>Stride of the source image in Y dimension (in bytes) </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_offset_first_element_in_bytes</td><td>The offset of the first element in the source image </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">src_step_z</td><td>src_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">dst_ptr</td><td>Pointer to the destination tensor. Supported data types: F32 </td></tr>
+ <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>
+ <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>
+ <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>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">dst_stride_z</td><td>Stride of the destination tensor in Z dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">dst_step_z</td><td>dst_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <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>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_ptr</td><td>Pointer to the weights tensor. Supported data types: F32 </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_x</td><td>Stride of the weights tensor in X dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_x</td><td>weights_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_y</td><td>Stride of the weights tensor in Y dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_y</td><td>weights_stride_y * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_stride_z</td><td>Stride of the weights tensor in Z dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_step_z</td><td>weights_stride_z * number of elements along Y processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">weights_offset_first_element_in_bytes</td><td>The offset of the first element in the biases vector </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_ptr</td><td>(Optional) Pointer to the biases vector. Supported data types: F32 </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_stride_x</td><td>(Optional) Stride of the biases vector in X dimension (in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_step_x</td><td>(Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) </td></tr>
+ <tr><td class="paramdir">[in]</td><td class="paramname">biases_offset_first_element_in_bytes</td><td>(Optional) The offset of the first element in the biases vector </td></tr>
+ </table>
+ </dd>
+</dl>
+
+<p>Definition at line <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00360">360</a> of file <a class="el" href="depthwise__convolution_8cl_source.xhtml">depthwise_convolution.cl</a>.</p>
+
+<p>References <a class="el" href="fixed__point_8h_source.xhtml#l00098">CONVERT</a>, <a class="el" href="helpers_8h_source.xhtml#l00114">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00117">CONVERT_TO_TENSOR3D_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00096">CONVERT_TO_VECTOR_STRUCT</a>, <a class="el" href="helpers_8h_source.xhtml#l00099">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>, <a class="el" href="depthwise__convolution_8cl_source.xhtml#l00221">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>, <a class="el" href="convolution3x3_8cl_source.xhtml#l00027">DATA_TYPE</a>, <a class="el" href="helpers_8h_source.xhtml#l00066">IMAGE_DECLARATION</a>, <a class="el" href="helpers_8h_source.xhtml#l00301">offset()</a>, <a class="el" href="helpers_8h_source.xhtml#l00134">Vector::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00142">Image::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00151">Tensor3D::ptr</a>, <a class="el" href="helpers_8h_source.xhtml#l00074">TENSOR3D_DECLARATION</a>, <a class="el" href="helpers_8h_source.xhtml#l00060">VECTOR_DECLARATION</a>, and <a class="el" href="helpers_8h_source.xhtml#l00290">vector_offset()</a>.</p>
+<div class="fragment"><div class="line"><a name="l00369"></a><span class="lineno"> 369</span> {</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span>  <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#a6743f0a130e8311e6f5b1a23df102472">src</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(src);</div><div class="line"><a name="l00371"></a><span class="lineno"> 371</span>  <a class="code" href="struct_image.xhtml">Image</a> <a class="code" href="namespacearm__compute_1_1test_1_1validation.xhtml#adbf67dcee294e673cf796f1ed8aeb6a4">dst</a> = <a class="code" href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a>(dst);</div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>  <a class="code" href="struct_tensor3_d.xhtml">Tensor3D</a> weights = <a class="code" href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a>(weights);</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span> </div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span>  float2 pixels0 = 0.0f;</div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>  float2 pixels1 = 0.0f;</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span> </div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>  __global uchar *weights_addr = (__global uchar *)weights.<a class="code" href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>;</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>  __global uchar *src_addr = (__global uchar *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, 0, 0);</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span> </div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>  <span class="comment">// Load the weights</span></div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>  float3 weights_row0 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 0 * weights_stride_y));</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>  float3 weights_row1 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 1 * weights_stride_y));</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>  float3 weights_row2 = vload3(0, (__global <span class="keywordtype">float</span> *)(weights_addr + 2 * weights_stride_y));</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span> </div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>  <span class="comment">// Note: Since each work-item computes 4x2 elements, we need to load 5 rows from the input tensor</span></div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span>  float4 src00 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span>  float2 src01 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 0 * src_stride_y)); <span class="comment">// Row0</span></div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>  float4 src10 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>  float2 src11 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 1 * src_stride_y)); <span class="comment">// Row1</span></div><div class="line"><a name="l00390"></a><span class="lineno"> 390</span>  float4 src20 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00391"></a><span class="lineno"> 391</span>  float2 src21 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 2 * src_stride_y)); <span class="comment">// Row2</span></div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span>  float4 src30 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span>  float2 src31 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 3 * src_stride_y)); <span class="comment">// Row3</span></div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span>  float4 src40 = vload4(0, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row4</span></div><div class="line"><a name="l00395"></a><span class="lineno"> 395</span>  float2 src41 = vload2(2, (__global <span class="keywordtype">float</span> *)(src_addr + 4 * src_stride_y)); <span class="comment">// Row4</span></div><div class="line"><a name="l00396"></a><span class="lineno"> 396</span> </div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src00, src01, weights_row0);</div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src10, src11, weights_row1);</div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels0, src20, src21, weights_row2);</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src20, src21, weights_row0);</div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src30, src31, weights_row1);</div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span>  <a class="code" href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a>(pixels1, src40, src41, weights_row2);</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span> </div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span> <span class="preprocessor">#ifdef HAS_BIAS</span></div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>  <a class="code" href="struct_vector.xhtml">Vector</a> biases = <a class="code" href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a>(biases);</div><div class="line"><a name="l00406"></a><span class="lineno"> 406</span> </div><div class="line"><a name="l00407"></a><span class="lineno"> 407</span>  <span class="keywordtype">float</span> bias = *((__global <span class="keywordtype">float</span> *)(<a class="code" href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a>(&biases, get_global_id(2))));</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span> </div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>  pixels0 += (float2)bias;</div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span>  pixels1 += (float2)bias;</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span> <span class="preprocessor">#endif </span><span class="comment">/* defined(HAS_BIAS) */</span><span class="preprocessor"></span></div><div class="line"><a name="l00412"></a><span class="lineno"> 412</span> </div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>  vstore2(pixels0, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 0 * dst_stride_y));</div><div class="line"><a name="l00414"></a><span class="lineno"> 414</span>  vstore2(pixels1, 0, (__global <span class="keywordtype">float</span> *)(dst.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + 1 * dst_stride_y));</div><div class="line"><a name="l00415"></a><span class="lineno"> 415</span> }</div><div class="ttc" id="struct_vector_xhtml"><div class="ttname"><a href="struct_vector.xhtml">Vector</a></div><div class="ttdoc">Structure to hold Vector information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00132">helpers.h:132</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a31c8c760f08fb1a331b16b7c204321dc"><div class="ttname"><a href="helpers_8h.xhtml#a31c8c760f08fb1a331b16b7c204321dc">CONVERT_TO_TENSOR3D_STRUCT</a></div><div class="ttdeci">#define CONVERT_TO_TENSOR3D_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00117">helpers.h:117</a></div></div>
+<div class="ttc" id="struct_tensor3_d_xhtml"><div class="ttname"><a href="struct_tensor3_d.xhtml">Tensor3D</a></div><div class="ttdoc">Structure to hold 3D tensor information. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00149">helpers.h:149</a></div></div>
+<div class="ttc" id="depthwise__convolution_8cl_xhtml_a31cbdd6b28ee6331839432f69d87fdcb"><div class="ttname"><a href="depthwise__convolution_8cl.xhtml#a31cbdd6b28ee6331839432f69d87fdcb">CONVOLUTION1x3_BIFROST2X1_STRIDE2</a></div><div class="ttdeci">#define CONVOLUTION1x3_BIFROST2X1_STRIDE2(acc, src0, src1, weights_row0)</div><div class="ttdef"><b>Definition:</b> <a href="depthwise__convolution_8cl_source.xhtml#l00221">depthwise_convolution.cl:221</a></div></div>
+<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>
+<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>
+<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>
+<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>
+<div class="ttc" id="helpers_8h_xhtml_a64d779f80eeb923e0ab2313433f7b40b"><div class="ttname"><a href="helpers_8h.xhtml#a64d779f80eeb923e0ab2313433f7b40b">CONVERT_TO_VECTOR_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a7e4940407322d6f0ccb8b6b86b856019"><div class="ttname"><a href="helpers_8h.xhtml#a7e4940407322d6f0ccb8b6b86b856019">vector_offset</a></div><div class="ttdeci">__global const uchar * vector_offset(const Vector *vec, int x)</div><div class="ttdoc">Get the pointer position of a Vector. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00290">helpers.h:290</a></div></div>
+<div class="ttc" id="helpers_8h_xhtml_a541f8db866a0fa93ee67d58ea31a7d0c"><div class="ttname"><a href="helpers_8h.xhtml#a541f8db866a0fa93ee67d58ea31a7d0c">CONVERT_TENSOR3D_TO_IMAGE_STRUCT</a></div><div class="ttdeci">#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00114">helpers.h:114</a></div></div>
+<div class="ttc" id="struct_tensor3_d_xhtml_acf52c23cbd7424606c10a606524e3e32"><div class="ttname"><a href="struct_tensor3_d.xhtml#acf52c23cbd7424606c10a606524e3e32">Tensor3D::ptr</a></div><div class="ttdeci">__global uchar * ptr</div><div class="ttdoc">Pointer to the starting postion of the buffer. </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00151">helpers.h:151</a></div></div>
+<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 & src</div><div class="ttdef"><b>Definition:</b> <a href="_c_l_2_convolution_8cpp_source.xhtml#l00147">Convolution.cpp:147</a></div></div>
+</div><!-- fragment -->
+</div>
+</div>
</div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
<ul>
<li class="navelem"><a class="el" href="dir_68267d1309a1af8e8297ef4c3efbcdba.xhtml">src</a></li><li class="navelem"><a class="el" href="dir_aebb8dcc11953d78e620bbef0b9e2183.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_8c278f79c760e5c5fbd911f9870614c1.xhtml">CL</a></li><li class="navelem"><a class="el" href="dir_25885286e9dad4fa105b7b25a8031bbf.xhtml">cl_kernels</a></li><li class="navelem"><a class="el" href="depthwise__convolution_8cl.xhtml">depthwise_convolution.cl</a></li>
- <li class="footer">Generated on Wed Jan 24 2018 14:30:43 for Compute Library by
+ <li class="footer">Generated on Thu Feb 22 2018 15:45:22 for Compute Library by
<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
</ul>