1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
|
<!-- HTML header for doxygen 1.8.9.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
<title>ARM Compute Library: src/core/CL/cl_kernels/canny.cl Source File</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtreedata.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
$(document).ready(initResizable);
$(window).load(resizeHeight);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td style="padding-left: 0.5em;">
<div id="projectname">ARM Compute Library
 <span id="projectnumber">17.05</span>
</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.xhtml"><span>Main Page</span></a></li>
<li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
<li><a href="annotated.xhtml"><span>Data Structures</span></a></li>
<li class="current"><a href="files.xhtml"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.xhtml"><span>File List</span></a></li>
<li><a href="globals.xhtml"><span>Globals</span></a></li>
</ul>
</div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
<div id="nav-tree">
<div id="nav-tree-contents">
<div id="nav-sync" class="sync"></div>
</div>
</div>
<div id="splitbar" style="-moz-user-select:none;"
class="ui-resizable-handle">
</div>
</div>
<script type="text/javascript">
$(document).ready(function(){initNavTree('canny_8cl_source.xhtml','');});
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div class="header">
<div class="headertitle">
<div class="title">canny.cl</div> </div>
</div><!--header-->
<div class="contents">
<a href="canny_8cl.xhtml">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno"> 1</span> <span class="comment">/*</span></div><div class="line"><a name="l00002"></a><span class="lineno"> 2</span> <span class="comment"> * Copyright (c) 2017 ARM Limited.</span></div><div class="line"><a name="l00003"></a><span class="lineno"> 3</span> <span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno"> 4</span> <span class="comment"> * SPDX-License-Identifier: MIT</span></div><div class="line"><a name="l00005"></a><span class="lineno"> 5</span> <span class="comment"> *</span></div><div class="line"><a name="l00006"></a><span class="lineno"> 6</span> <span class="comment"> * Permission is hereby granted, free of charge, to any person obtaining a copy</span></div><div class="line"><a name="l00007"></a><span class="lineno"> 7</span> <span class="comment"> * of this software and associated documentation files (the "Software"), to</span></div><div class="line"><a name="l00008"></a><span class="lineno"> 8</span> <span class="comment"> * deal in the Software without restriction, including without limitation the</span></div><div class="line"><a name="l00009"></a><span class="lineno"> 9</span> <span class="comment"> * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or</span></div><div class="line"><a name="l00010"></a><span class="lineno"> 10</span> <span class="comment"> * sell copies of the Software, and to permit persons to whom the Software is</span></div><div class="line"><a name="l00011"></a><span class="lineno"> 11</span> <span class="comment"> * furnished to do so, subject to the following conditions:</span></div><div class="line"><a name="l00012"></a><span class="lineno"> 12</span> <span class="comment"> *</span></div><div class="line"><a name="l00013"></a><span class="lineno"> 13</span> <span class="comment"> * The above copyright notice and this permission notice shall be included in all</span></div><div class="line"><a name="l00014"></a><span class="lineno"> 14</span> <span class="comment"> * copies or substantial portions of the Software.</span></div><div class="line"><a name="l00015"></a><span class="lineno"> 15</span> <span class="comment"> *</span></div><div class="line"><a name="l00016"></a><span class="lineno"> 16</span> <span class="comment"> * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</span></div><div class="line"><a name="l00017"></a><span class="lineno"> 17</span> <span class="comment"> * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,</span></div><div class="line"><a name="l00018"></a><span class="lineno"> 18</span> <span class="comment"> * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE</span></div><div class="line"><a name="l00019"></a><span class="lineno"> 19</span> <span class="comment"> * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER</span></div><div class="line"><a name="l00020"></a><span class="lineno"> 20</span> <span class="comment"> * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,</span></div><div class="line"><a name="l00021"></a><span class="lineno"> 21</span> <span class="comment"> * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE</span></div><div class="line"><a name="l00022"></a><span class="lineno"> 22</span> <span class="comment"> * SOFTWARE.</span></div><div class="line"><a name="l00023"></a><span class="lineno"> 23</span> <span class="comment"> */</span></div><div class="line"><a name="l00024"></a><span class="lineno"> 24</span> <span class="preprocessor">#include "<a class="code" href="helpers_8h.xhtml">helpers.h</a>"</span></div><div class="line"><a name="l00025"></a><span class="lineno"> 25</span> </div><div class="line"><a name="l00057"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#a92f95668118a1bc6ef0eed965dfd570f"> 57</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="canny_8cl.xhtml#a92f95668118a1bc6ef0eed965dfd570f">combine_gradients_L1</a>(</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src1),</div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src2),</div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(grad),</div><div class="line"><a name="l00061"></a><span class="lineno"> 61</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(angle))</div><div class="line"><a name="l00062"></a><span class="lineno"> 62</span> {</div><div class="line"><a name="l00063"></a><span class="lineno"> 63</span>  <span class="comment">// Construct images</span></div><div class="line"><a name="l00064"></a><span class="lineno"> 64</span>  <a class="code" href="struct_image.xhtml">Image</a> src1 = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src1);</div><div class="line"><a name="l00065"></a><span class="lineno"> 65</span>  <a class="code" href="struct_image.xhtml">Image</a> src2 = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src2);</div><div class="line"><a name="l00066"></a><span class="lineno"> 66</span>  <a class="code" href="struct_image.xhtml">Image</a> grad = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(grad);</div><div class="line"><a name="l00067"></a><span class="lineno"> 67</span>  <a class="code" href="struct_image.xhtml">Image</a> angle = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(angle);</div><div class="line"><a name="l00068"></a><span class="lineno"> 68</span> </div><div class="line"><a name="l00069"></a><span class="lineno"> 69</span>  <span class="comment">// Load sobel horizontal and vertical values</span></div><div class="line"><a name="l00070"></a><span class="lineno"> 70</span>  <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE_IN, 4)</div><div class="line"><a name="l00071"></a><span class="lineno"> 71</span>  h = vload4(0, (__global DATA_TYPE_IN *)src1.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00072"></a><span class="lineno"> 72</span>  <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE_IN, 4)</div><div class="line"><a name="l00073"></a><span class="lineno"> 73</span>  v = vload4(0, (__global DATA_TYPE_IN *)src2.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00074"></a><span class="lineno"> 74</span> </div><div class="line"><a name="l00075"></a><span class="lineno"> 75</span>  <span class="comment">/* Calculate the gradient, using level 1 normalisation method */</span></div><div class="line"><a name="l00076"></a><span class="lineno"> 76</span>  <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a>, 4)</div><div class="line"><a name="l00077"></a><span class="lineno"> 77</span>  m = <a class="code" href="helpers_8h.xhtml#a23fb01b6f3453cc0e48a026fd44f6acd">CONVERT_SAT</a>((abs(h) + abs(v)), <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a>, 4));</div><div class="line"><a name="l00078"></a><span class="lineno"> 78</span> </div><div class="line"><a name="l00079"></a><span class="lineno"> 79</span>  <span class="comment">/* Calculate the angle */</span></div><div class="line"><a name="l00080"></a><span class="lineno"> 80</span>  float4 p = atan2pi(convert_float4(v), convert_float4(h));</div><div class="line"><a name="l00081"></a><span class="lineno"> 81</span> </div><div class="line"><a name="l00082"></a><span class="lineno"> 82</span>  <span class="comment">/* Remap angle to range [0, 256) */</span></div><div class="line"><a name="l00083"></a><span class="lineno"> 83</span>  p = select(p, p + 2, p < 0.0f) * 128.0f;</div><div class="line"><a name="l00084"></a><span class="lineno"> 84</span> </div><div class="line"><a name="l00085"></a><span class="lineno"> 85</span>  <span class="comment">/* Store results */</span></div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span>  vstore4(m, 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a> *)grad.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00087"></a><span class="lineno"> 87</span>  vstore4(convert_uchar4_sat_rte(p), 0, angle.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span> }</div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span> </div><div class="line"><a name="l00121"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#a7046357526e2f95a34e3ffa217752d9f"> 121</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="canny_8cl.xhtml#a7046357526e2f95a34e3ffa217752d9f">combine_gradients_L2</a>(</div><div class="line"><a name="l00122"></a><span class="lineno"> 122</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src1),</div><div class="line"><a name="l00123"></a><span class="lineno"> 123</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src2),</div><div class="line"><a name="l00124"></a><span class="lineno"> 124</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(grad),</div><div class="line"><a name="l00125"></a><span class="lineno"> 125</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(angle))</div><div class="line"><a name="l00126"></a><span class="lineno"> 126</span> {</div><div class="line"><a name="l00127"></a><span class="lineno"> 127</span>  <span class="comment">// Construct images</span></div><div class="line"><a name="l00128"></a><span class="lineno"> 128</span>  <a class="code" href="struct_image.xhtml">Image</a> src1 = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src1);</div><div class="line"><a name="l00129"></a><span class="lineno"> 129</span>  <a class="code" href="struct_image.xhtml">Image</a> src2 = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(src2);</div><div class="line"><a name="l00130"></a><span class="lineno"> 130</span>  <a class="code" href="struct_image.xhtml">Image</a> grad = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(grad);</div><div class="line"><a name="l00131"></a><span class="lineno"> 131</span>  <a class="code" href="struct_image.xhtml">Image</a> angle = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(angle);</div><div class="line"><a name="l00132"></a><span class="lineno"> 132</span> </div><div class="line"><a name="l00133"></a><span class="lineno"> 133</span>  <span class="comment">// Load sobel horizontal and vertical values</span></div><div class="line"><a name="l00134"></a><span class="lineno"> 134</span>  float4 h = convert_float4(vload4(0, (__global DATA_TYPE_IN *)src1.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>));</div><div class="line"><a name="l00135"></a><span class="lineno"> 135</span>  float4 v = convert_float4(vload4(0, (__global DATA_TYPE_IN *)src2.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>));</div><div class="line"><a name="l00136"></a><span class="lineno"> 136</span> </div><div class="line"><a name="l00137"></a><span class="lineno"> 137</span>  <span class="comment">/* Calculate the gradient, using level 2 normalisation method */</span></div><div class="line"><a name="l00138"></a><span class="lineno"> 138</span>  float4 m = sqrt(h * h + v * v);</div><div class="line"><a name="l00139"></a><span class="lineno"> 139</span> </div><div class="line"><a name="l00140"></a><span class="lineno"> 140</span>  <span class="comment">/* Calculate the angle */</span></div><div class="line"><a name="l00141"></a><span class="lineno"> 141</span>  float4 p = atan2pi(v, h);</div><div class="line"><a name="l00142"></a><span class="lineno"> 142</span> </div><div class="line"><a name="l00143"></a><span class="lineno"> 143</span>  <span class="comment">/* Remap angle to range [0, 256) */</span></div><div class="line"><a name="l00144"></a><span class="lineno"> 144</span>  p = select(p, p + 2, p < 0.0f) * 128.0f;</div><div class="line"><a name="l00145"></a><span class="lineno"> 145</span> </div><div class="line"><a name="l00146"></a><span class="lineno"> 146</span>  <span class="comment">/* Store results */</span></div><div class="line"><a name="l00147"></a><span class="lineno"> 147</span>  vstore4(<a class="code" href="helpers_8h.xhtml#a28b2b123792a5b11a32cb05a6d327437">CONVERT_SAT_ROUND</a>(m, <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(<a class="code" href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a>, 4), rte), 0, (__global <a class="code" href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a> *)grad.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00148"></a><span class="lineno"> 148</span>  vstore4(convert_uchar4_sat_rte(p), 0, angle.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00149"></a><span class="lineno"> 149</span> }</div><div class="line"><a name="l00150"></a><span class="lineno"> 150</span> </div><div class="line"><a name="l00153"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#a65e45b2f8637f10dd14a427f37c2cd4c"> 153</a></span> __constant short4 <a class="code" href="canny_8cl.xhtml#a65e45b2f8637f10dd14a427f37c2cd4c">neighbours_coords</a>[] =</div><div class="line"><a name="l00154"></a><span class="lineno"> 154</span> {</div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span>  { -1, 0, 1, 0 }, <span class="comment">// 0</span></div><div class="line"><a name="l00156"></a><span class="lineno"> 156</span>  { -1, 1, 1, -1 }, <span class="comment">// 45</span></div><div class="line"><a name="l00157"></a><span class="lineno"> 157</span>  { 0, 1, 0, -1 }, <span class="comment">// 90</span></div><div class="line"><a name="l00158"></a><span class="lineno"> 158</span>  { 1, 1, -1, -1 }, <span class="comment">// 135</span></div><div class="line"><a name="l00159"></a><span class="lineno"> 159</span>  { 1, 0, -1, 0 }, <span class="comment">// 180</span></div><div class="line"><a name="l00160"></a><span class="lineno"> 160</span>  { 1, -1, -1, 1 }, <span class="comment">// 225</span></div><div class="line"><a name="l00161"></a><span class="lineno"> 161</span>  { 0, 1, 0, -1 }, <span class="comment">// 270</span></div><div class="line"><a name="l00162"></a><span class="lineno"> 162</span>  { -1, -1, 1, 1 }, <span class="comment">// 315</span></div><div class="line"><a name="l00163"></a><span class="lineno"> 163</span>  { -1, 0, 1, 0 }, <span class="comment">// 360</span></div><div class="line"><a name="l00164"></a><span class="lineno"> 164</span> };</div><div class="line"><a name="l00165"></a><span class="lineno"> 165</span> </div><div class="line"><a name="l00191"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#a8c233c75ee1010bcec8601d559f4be68"> 191</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="canny_8cl.xhtml#a8c233c75ee1010bcec8601d559f4be68">suppress_non_maximum</a>(</div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(grad),</div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(angle),</div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(non_max),</div><div class="line"><a name="l00195"></a><span class="lineno"> 195</span>  uint lower_thr)</div><div class="line"><a name="l00196"></a><span class="lineno"> 196</span> {</div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>  <span class="comment">// Construct images</span></div><div class="line"><a name="l00198"></a><span class="lineno"> 198</span>  <a class="code" href="struct_image.xhtml">Image</a> grad = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(grad);</div><div class="line"><a name="l00199"></a><span class="lineno"> 199</span>  <a class="code" href="struct_image.xhtml">Image</a> angle = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(angle);</div><div class="line"><a name="l00200"></a><span class="lineno"> 200</span>  <a class="code" href="struct_image.xhtml">Image</a> non_max = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(non_max);</div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span> </div><div class="line"><a name="l00202"></a><span class="lineno"> 202</span>  <span class="comment">// Get gradient and angle</span></div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>  DATA_TYPE_IN gradient = *((__global DATA_TYPE_IN *)grad.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</span>  uchar an = convert_ushort(*angle.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00205"></a><span class="lineno"> 205</span> </div><div class="line"><a name="l00206"></a><span class="lineno"> 206</span>  <span class="keywordflow">if</span>(gradient <= lower_thr)</div><div class="line"><a name="l00207"></a><span class="lineno"> 207</span>  {</div><div class="line"><a name="l00208"></a><span class="lineno"> 208</span>  <span class="keywordflow">return</span>;</div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span>  }</div><div class="line"><a name="l00210"></a><span class="lineno"> 210</span> </div><div class="line"><a name="l00211"></a><span class="lineno"> 211</span>  <span class="comment">// Divide the whole round into 8 directions</span></div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span>  uchar ang = 127 - an;</div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>  <a class="code" href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a> q_an = (ang + 16) >> 5;</div><div class="line"><a name="l00214"></a><span class="lineno"> 214</span> </div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>  <span class="comment">// Find the two pixels in the perpendicular direction</span></div><div class="line"><a name="l00216"></a><span class="lineno"> 216</span>  short2 x_p = neighbours_coords[q_an].s02;</div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>  short2 y_p = neighbours_coords[q_an].s13;</div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>  DATA_TYPE_IN g1 = *((global DATA_TYPE_IN *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&grad, x_p.x, y_p.x));</div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span>  DATA_TYPE_IN g2 = *((global DATA_TYPE_IN *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&grad, x_p.y, y_p.y));</div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span> </div><div class="line"><a name="l00221"></a><span class="lineno"> 221</span>  <span class="keywordflow">if</span>((gradient > g1) && (gradient > g2))</div><div class="line"><a name="l00222"></a><span class="lineno"> 222</span>  {</div><div class="line"><a name="l00223"></a><span class="lineno"> 223</span>  *((global <a class="code" href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a> *)non_max.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>) = gradient;</div><div class="line"><a name="l00224"></a><span class="lineno"> 224</span>  }</div><div class="line"><a name="l00225"></a><span class="lineno"> 225</span> }</div><div class="line"><a name="l00226"></a><span class="lineno"> 226</span> </div><div class="line"><a name="l00227"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#ac243bfa96aa2c28014159ff098bd2324"> 227</a></span> <span class="preprocessor">#define EDGE 255</span></div><div class="line"><a name="l00228"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#a242e5ec23ae26cd2dfa5eab5ae68a0f1"> 228</a></span> <span class="preprocessor">#define hysteresis_local_stack_L1 8 // The size of level 1 stack. This has to agree with the host side</span></div><div class="line"><a name="l00229"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#aa55093fb165fb76898efb3ba80f62bae"> 229</a></span> <span class="preprocessor">#define hysteresis_local_stack_L2 16 // The size of level 2 stack, adjust this can impact the match rate with VX implementation</span></div><div class="line"><a name="l00230"></a><span class="lineno"> 230</span> </div><div class="line"><a name="l00242"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650"> 242</a></span> <span class="preprocessor">#define check_pixel(early_test, x_pos, y_pos, x_cur, y_cur) \</span></div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span> <span class="preprocessor"> { \</span></div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span> <span class="preprocessor"> if(!early_test) \</span></div><div class="line"><a name="l00245"></a><span class="lineno"> 245</span> <span class="preprocessor"> { \</span></div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span> <span class="preprocessor"> </span><span class="comment">/* Number of elements in the local stack 1, points to next available entry */</span><span class="preprocessor"> \</span></div><div class="line"><a name="l00247"></a><span class="lineno"> 247</span> <span class="preprocessor"> c = *((__global char *)offset(&l1_stack_counter, x_cur, y_cur)); \</span></div><div class="line"><a name="l00248"></a><span class="lineno"> 248</span> <span class="preprocessor"> \</span></div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span> <span class="preprocessor"> if(c > (hysteresis_local_stack_L1 - 1)) </span><span class="comment">/* Stack level 1 is full */</span><span class="preprocessor"> \</span></div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span> <span class="preprocessor"> goto pop_stack; \</span></div><div class="line"><a name="l00251"></a><span class="lineno"> 251</span> <span class="preprocessor"> \</span></div><div class="line"><a name="l00252"></a><span class="lineno"> 252</span> <span class="preprocessor"> </span><span class="comment">/* The pixel that has already been recorded is ignored */</span><span class="preprocessor"> \</span></div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span> <span class="preprocessor"> if(!atomic_or((__global uint *)offset(&recorded, x_pos, y_pos), 1)) \</span></div><div class="line"><a name="l00254"></a><span class="lineno"> 254</span> <span class="preprocessor"> { \</span></div><div class="line"><a name="l00255"></a><span class="lineno"> 255</span> <span class="preprocessor"> l1_ptr[c] = (short2)(x_pos, y_pos); \</span></div><div class="line"><a name="l00256"></a><span class="lineno"> 256</span> <span class="preprocessor"> *((__global char *)offset(&l1_stack_counter, x_cur, y_cur)) += 1; \</span></div><div class="line"><a name="l00257"></a><span class="lineno"> 257</span> <span class="preprocessor"> } \</span></div><div class="line"><a name="l00258"></a><span class="lineno"> 258</span> <span class="preprocessor"> } \</span></div><div class="line"><a name="l00259"></a><span class="lineno"> 259</span> <span class="preprocessor"> }</span></div><div class="line"><a name="l00260"></a><span class="lineno"> 260</span> </div><div class="line"><a name="l00306"></a><span class="lineno"><a class="line" href="canny_8cl.xhtml#acd62ae1c9f3d7a1c7e49499d308d1904"> 306</a></span> kernel <span class="keywordtype">void</span> <a class="code" href="canny_8cl.xhtml#acd62ae1c9f3d7a1c7e49499d308d1904">hysteresis</a>(</div><div class="line"><a name="l00307"></a><span class="lineno"> 307</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(src),</div><div class="line"><a name="l00308"></a><span class="lineno"> 308</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(out),</div><div class="line"><a name="l00309"></a><span class="lineno"> 309</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(visited),</div><div class="line"><a name="l00310"></a><span class="lineno"> 310</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(recorded),</div><div class="line"><a name="l00311"></a><span class="lineno"> 311</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(l1_stack),</div><div class="line"><a name="l00312"></a><span class="lineno"> 312</span>  <a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(l1_stack_counter),</div><div class="line"><a name="l00313"></a><span class="lineno"> 313</span>  uint low_thr,</div><div class="line"><a name="l00314"></a><span class="lineno"> 314</span>  uint up_thr,</div><div class="line"><a name="l00315"></a><span class="lineno"> 315</span>  <span class="keywordtype">int</span> width,</div><div class="line"><a name="l00316"></a><span class="lineno"> 316</span>  <span class="keywordtype">int</span> height)</div><div class="line"><a name="l00317"></a><span class="lineno"> 317</span> {</div><div class="line"><a name="l00318"></a><span class="lineno"> 318</span>  <span class="comment">// Create images</span></div><div class="line"><a name="l00319"></a><span class="lineno"> 319</span>  <a class="code" href="struct_image.xhtml">Image</a> src = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(src);</div><div class="line"><a name="l00320"></a><span class="lineno"> 320</span>  <a class="code" href="struct_image.xhtml">Image</a> out = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(out);</div><div class="line"><a name="l00321"></a><span class="lineno"> 321</span>  <a class="code" href="struct_image.xhtml">Image</a> visited = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(visited);</div><div class="line"><a name="l00322"></a><span class="lineno"> 322</span>  <a class="code" href="struct_image.xhtml">Image</a> recorded = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(recorded);</div><div class="line"><a name="l00323"></a><span class="lineno"> 323</span>  <a class="code" href="struct_image.xhtml">Image</a> l1_stack = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(l1_stack);</div><div class="line"><a name="l00324"></a><span class="lineno"> 324</span>  <a class="code" href="struct_image.xhtml">Image</a> l1_stack_counter = <a class="code" href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a>(l1_stack_counter);</div><div class="line"><a name="l00325"></a><span class="lineno"> 325</span> </div><div class="line"><a name="l00326"></a><span class="lineno"> 326</span>  <span class="comment">// Index</span></div><div class="line"><a name="l00327"></a><span class="lineno"> 327</span>  <span class="keywordtype">int</span> x = get_global_id(0);</div><div class="line"><a name="l00328"></a><span class="lineno"> 328</span>  <span class="keywordtype">int</span> y = get_global_id(1);</div><div class="line"><a name="l00329"></a><span class="lineno"> 329</span> </div><div class="line"><a name="l00330"></a><span class="lineno"> 330</span>  <span class="comment">// Load value</span></div><div class="line"><a name="l00331"></a><span class="lineno"> 331</span>  DATA_TYPE_IN val = *((__global DATA_TYPE_IN *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, x, y));</div><div class="line"><a name="l00332"></a><span class="lineno"> 332</span> </div><div class="line"><a name="l00333"></a><span class="lineno"> 333</span>  <span class="comment">// If less than upper threshold set to NO_EDGE and return</span></div><div class="line"><a name="l00334"></a><span class="lineno"> 334</span>  <span class="keywordflow">if</span>(val <= up_thr)</div><div class="line"><a name="l00335"></a><span class="lineno"> 335</span>  {</div><div class="line"><a name="l00336"></a><span class="lineno"> 336</span>  *<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&out, x, y) = 0;</div><div class="line"><a name="l00337"></a><span class="lineno"> 337</span>  <span class="keywordflow">return</span>;</div><div class="line"><a name="l00338"></a><span class="lineno"> 338</span>  }</div><div class="line"><a name="l00339"></a><span class="lineno"> 339</span> </div><div class="line"><a name="l00340"></a><span class="lineno"> 340</span>  <span class="comment">// Init local stack 2</span></div><div class="line"><a name="l00341"></a><span class="lineno"> 341</span>  short2 stack_L2[<a class="code" href="canny_8cl.xhtml#aa55093fb165fb76898efb3ba80f62bae">hysteresis_local_stack_L2</a>] = { 0 };</div><div class="line"><a name="l00342"></a><span class="lineno"> 342</span>  <span class="keywordtype">int</span> L2_counter = 0;</div><div class="line"><a name="l00343"></a><span class="lineno"> 343</span> </div><div class="line"><a name="l00344"></a><span class="lineno"> 344</span>  <span class="comment">// Perform recursive hysteresis</span></div><div class="line"><a name="l00345"></a><span class="lineno"> 345</span>  <span class="keywordflow">while</span>(<span class="keyword">true</span>)</div><div class="line"><a name="l00346"></a><span class="lineno"> 346</span>  {</div><div class="line"><a name="l00347"></a><span class="lineno"> 347</span>  <span class="comment">// Get L1 stack pointer</span></div><div class="line"><a name="l00348"></a><span class="lineno"> 348</span>  __global short2 *l1_ptr = (__global short2 *)(l1_stack.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a> + y * l1_stack.<a class="code" href="struct_image.xhtml#a4f0b90c9ecd6e57ceb3f37332fefe8f1">stride_y</a> + x * <a class="code" href="canny_8cl.xhtml#a242e5ec23ae26cd2dfa5eab5ae68a0f1">hysteresis_local_stack_L1</a> * l1_stack.<a class="code" href="struct_image.xhtml#ae01febbfd0689ef709f3ff6fdd2abc7e">stride_x</a>);</div><div class="line"><a name="l00349"></a><span class="lineno"> 349</span> </div><div class="line"><a name="l00350"></a><span class="lineno"> 350</span>  <span class="comment">// If the pixel has already been visited, proceed with the items in the stack instead</span></div><div class="line"><a name="l00351"></a><span class="lineno"> 351</span>  <span class="keywordflow">if</span>(atomic_or((__global uint *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&visited, x, y), 1) != 0)</div><div class="line"><a name="l00352"></a><span class="lineno"> 352</span>  {</div><div class="line"><a name="l00353"></a><span class="lineno"> 353</span>  <span class="keywordflow">goto</span> pop_stack;</div><div class="line"><a name="l00354"></a><span class="lineno"> 354</span>  }</div><div class="line"><a name="l00355"></a><span class="lineno"> 355</span> </div><div class="line"><a name="l00356"></a><span class="lineno"> 356</span>  <span class="comment">// Set strong edge</span></div><div class="line"><a name="l00357"></a><span class="lineno"> 357</span>  *<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&out, x, y) = <a class="code" href="canny_8cl.xhtml#ac243bfa96aa2c28014159ff098bd2324">EDGE</a>;</div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span> </div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span>  <span class="comment">// If it is the top of stack l2, we don't need check the surrounding pixels</span></div><div class="line"><a name="l00360"></a><span class="lineno"> 360</span>  <span class="keywordflow">if</span>(L2_counter > (<a class="code" href="canny_8cl.xhtml#aa55093fb165fb76898efb3ba80f62bae">hysteresis_local_stack_L2</a> - 1))</div><div class="line"><a name="l00361"></a><span class="lineno"> 361</span>  {</div><div class="line"><a name="l00362"></a><span class="lineno"> 362</span>  <span class="keywordflow">goto</span> pop_stack2;</div><div class="line"><a name="l00363"></a><span class="lineno"> 363</span>  }</div><div class="line"><a name="l00364"></a><span class="lineno"> 364</span> </div><div class="line"><a name="l00365"></a><span class="lineno"> 365</span>  <span class="comment">// Points to the start of the local stack;</span></div><div class="line"><a name="l00366"></a><span class="lineno"> 366</span>  <span class="keywordtype">char</span> c;</div><div class="line"><a name="l00367"></a><span class="lineno"> 367</span> </div><div class="line"><a name="l00368"></a><span class="lineno"> 368</span>  <a class="code" href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a>(DATA_TYPE_IN, 4)</div><div class="line"><a name="l00369"></a><span class="lineno"> 369</span>  x_tmp;</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span>  uint4 v_tmp;</div><div class="line"><a name="l00371"></a><span class="lineno"> 371</span> </div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>  <span class="comment">// Get direction pixel indices</span></div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span>  <span class="keywordtype">int</span> N = max(y - 1, 0), S = min(y + 1, height - 2), W = max(x - 1, 0), E = min(x + 1, width - 2);</div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span> </div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>  <span class="comment">// Check 8 pixels around for week edges where low_thr < val <= up_thr</span></div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>  x_tmp = vload4(0, (__global DATA_TYPE_IN *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, W, N));</div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>  v_tmp = vload4(0, (__global uint *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&visited, W, N));</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, N, x, y); <span class="comment">// NW</span></div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s1 <= low_thr) || v_tmp.s1 || (x_tmp.s1 > up_thr)), x, N, x, y); <span class="comment">// N</span></div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, N, x, y); <span class="comment">// NE</span></div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span> </div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>  x_tmp = vload4(0, (__global DATA_TYPE_IN *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, W, y));</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>  v_tmp = vload4(0, (__global uint *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&visited, W, y));</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, y, x, y); <span class="comment">// W</span></div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, y, x, y); <span class="comment">// E</span></div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span> </div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span>  x_tmp = vload4(0, (__global DATA_TYPE_IN *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&src, W, S));</div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>  v_tmp = vload4(0, (__global uint *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&visited, W, S));</div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, S, x, y); <span class="comment">// SW</span></div><div class="line"><a name="l00390"></a><span class="lineno"> 390</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s1 <= low_thr) || v_tmp.s1 || (x_tmp.s1 > up_thr)), x, S, x, y); <span class="comment">// S</span></div><div class="line"><a name="l00391"></a><span class="lineno"> 391</span>  <a class="code" href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a>(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, S, x, y); <span class="comment">// SE</span></div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span> </div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span> <span class="preprocessor">#undef check_pixel</span></div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span> </div><div class="line"><a name="l00395"></a><span class="lineno"> 395</span> pop_stack:</div><div class="line"><a name="l00396"></a><span class="lineno"> 396</span>  c = *((__global <span class="keywordtype">char</span> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&l1_stack_counter, x, y));</div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span> </div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>  <span class="keywordflow">if</span>(c >= 1)</div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>  {</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span>  *((__global <span class="keywordtype">char</span> *)<a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>(&l1_stack_counter, x, y)) -= 1;</div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>  int2 l_c = convert_int2(l1_ptr[c - 1]);</div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span> </div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>  <span class="comment">// Push the current position into level 2 stack</span></div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>  stack_L2[L2_counter].x = x;</div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>  stack_L2[L2_counter].y = y;</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>  x = l_c.x;</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span>  y = l_c.y;</div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span> </div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span>  L2_counter++;</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span> </div><div class="line"><a name="l00412"></a><span class="lineno"> 412</span>  <span class="keywordflow">continue</span>;</div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>  }</div><div class="line"><a name="l00414"></a><span class="lineno"> 414</span> </div><div class="line"><a name="l00415"></a><span class="lineno"> 415</span>  <span class="keywordflow">if</span>(L2_counter > 0)</div><div class="line"><a name="l00416"></a><span class="lineno"> 416</span>  {</div><div class="line"><a name="l00417"></a><span class="lineno"> 417</span>  <span class="keywordflow">goto</span> pop_stack2;</div><div class="line"><a name="l00418"></a><span class="lineno"> 418</span>  }</div><div class="line"><a name="l00419"></a><span class="lineno"> 419</span>  <span class="keywordflow">else</span></div><div class="line"><a name="l00420"></a><span class="lineno"> 420</span>  {</div><div class="line"><a name="l00421"></a><span class="lineno"> 421</span>  <span class="keywordflow">return</span>;</div><div class="line"><a name="l00422"></a><span class="lineno"> 422</span>  }</div><div class="line"><a name="l00423"></a><span class="lineno"> 423</span> </div><div class="line"><a name="l00424"></a><span class="lineno"> 424</span> pop_stack2:</div><div class="line"><a name="l00425"></a><span class="lineno"> 425</span>  L2_counter--;</div><div class="line"><a name="l00426"></a><span class="lineno"> 426</span>  x = stack_L2[L2_counter].x;</div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>  y = stack_L2[L2_counter].y;</div><div class="line"><a name="l00428"></a><span class="lineno"> 428</span>  };</div><div class="line"><a name="l00429"></a><span class="lineno"> 429</span> }</div><div class="ttc" id="canny_8cl_xhtml_a242e5ec23ae26cd2dfa5eab5ae68a0f1"><div class="ttname"><a href="canny_8cl.xhtml#a242e5ec23ae26cd2dfa5eab5ae68a0f1">hysteresis_local_stack_L1</a></div><div class="ttdeci">#define hysteresis_local_stack_L1</div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00228">canny.cl:228</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_a8c233c75ee1010bcec8601d559f4be68"><div class="ttname"><a href="canny_8cl.xhtml#a8c233c75ee1010bcec8601d559f4be68">suppress_non_maximum</a></div><div class="ttdeci">__kernel void suppress_non_maximum(__global uchar *grad_ptr, uint grad_stride_x, uint grad_step_x, uint grad_stride_y, uint grad_step_y, uint grad_offset_first_element_in_bytes, __global uchar *angle_ptr, uint angle_stride_x, uint angle_step_x, uint angle_stride_y, uint angle_step_y, uint angle_offset_first_element_in_bytes, __global uchar *non_max_ptr, uint non_max_stride_x, uint non_max_step_x, uint non_max_stride_y, uint non_max_step_y, uint non_max_offset_first_element_in_bytes, uint lower_thr)</div><div class="ttdoc">Perform non maximum suppression. </div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00191">canny.cl:191</a></div></div>
<div class="ttc" id="helpers_8h_xhtml_a4334a4a76f8e9628c0fb9e1acf616e2a"><div class="ttname"><a href="helpers_8h.xhtml#a4334a4a76f8e9628c0fb9e1acf616e2a">CONVERT_TO_IMAGE_STRUCT_NO_STEP</a></div><div class="ttdeci">#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00076">helpers.h:76</a></div></div>
<div class="ttc" id="helpers_8h_xhtml_a22f42fcf2077d951271df83b55c1a71a"><div class="ttname"><a href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a></div><div class="ttdeci">#define IMAGE_DECLARATION(name)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00049">helpers.h:49</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#l00201">helpers.h:201</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_a65e45b2f8637f10dd14a427f37c2cd4c"><div class="ttname"><a href="canny_8cl.xhtml#a65e45b2f8637f10dd14a427f37c2cd4c">neighbours_coords</a></div><div class="ttdeci">__constant short4 neighbours_coords[]</div><div class="ttdoc">Array that holds the relative coordinates offset for the neighbouring pixels. </div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00153">canny.cl:153</a></div></div>
<div class="ttc" id="helpers_8h_xhtml"><div class="ttname"><a href="helpers_8h.xhtml">helpers.h</a></div></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#l00073">helpers.h:73</a></div></div>
<div class="ttc" id="helpers_8h_xhtml_a36f754c05b6fddf6df0d8d0a74f8159f"><div class="ttname"><a href="helpers_8h.xhtml#a36f754c05b6fddf6df0d8d0a74f8159f">VEC_DATA_TYPE</a></div><div class="ttdeci">#define VEC_DATA_TYPE(type, size)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00032">helpers.h:32</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_ac243bfa96aa2c28014159ff098bd2324"><div class="ttname"><a href="canny_8cl.xhtml#ac243bfa96aa2c28014159ff098bd2324">EDGE</a></div><div class="ttdeci">#define EDGE</div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00227">canny.cl:227</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#l00095">helpers.h:95</a></div></div>
<div class="ttc" id="helpers_8h_xhtml_a23fb01b6f3453cc0e48a026fd44f6acd"><div class="ttname"><a href="helpers_8h.xhtml#a23fb01b6f3453cc0e48a026fd44f6acd">CONVERT_SAT</a></div><div class="ttdeci">#define CONVERT_SAT(x, type)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00038">helpers.h:38</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#l00097">helpers.h:97</a></div></div>
<div class="ttc" id="helpers_8h_xhtml_a28b2b123792a5b11a32cb05a6d327437"><div class="ttname"><a href="helpers_8h.xhtml#a28b2b123792a5b11a32cb05a6d327437">CONVERT_SAT_ROUND</a></div><div class="ttdeci">#define CONVERT_SAT_ROUND(x, type, round)</div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00041">helpers.h:41</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_a64ee229d1bcee88c8017356d5d485650"><div class="ttname"><a href="canny_8cl.xhtml#a64ee229d1bcee88c8017356d5d485650">check_pixel</a></div><div class="ttdeci">#define check_pixel(early_test, x_pos, y_pos, x_cur, y_cur)</div><div class="ttdoc">Check whether pixel is valid. </div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00242">canny.cl:242</a></div></div>
<div class="ttc" id="convolution3x3_8cl_xhtml_ac06f3e24d3fffd3c465d8b2a6e7c985e"><div class="ttname"><a href="convolution3x3_8cl.xhtml#ac06f3e24d3fffd3c465d8b2a6e7c985e">DATA_TYPE_OUT</a></div><div class="ttdeci">#define DATA_TYPE_OUT</div><div class="ttdef"><b>Definition:</b> <a href="convolution3x3_8cl_source.xhtml#l00031">convolution3x3.cl:31</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_aa55093fb165fb76898efb3ba80f62bae"><div class="ttname"><a href="canny_8cl.xhtml#aa55093fb165fb76898efb3ba80f62bae">hysteresis_local_stack_L2</a></div><div class="ttdeci">#define hysteresis_local_stack_L2</div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00229">canny.cl:229</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_a7046357526e2f95a34e3ffa217752d9f"><div class="ttname"><a href="canny_8cl.xhtml#a7046357526e2f95a34e3ffa217752d9f">combine_gradients_L2</a></div><div class="ttdeci">__kernel void combine_gradients_L2(__global uchar *src1_ptr, uint src1_stride_x, uint src1_step_x, uint src1_stride_y, uint src1_step_y, uint src1_offset_first_element_in_bytes, __global uchar *src2_ptr, uint src2_stride_x, uint src2_step_x, uint src2_stride_y, uint src2_step_y, uint src2_offset_first_element_in_bytes, __global uchar *grad_ptr, uint grad_stride_x, uint grad_step_x, uint grad_stride_y, uint grad_step_y, uint grad_offset_first_element_in_bytes, __global uchar *angle_ptr, uint angle_stride_x, uint angle_step_x, uint angle_stride_y, uint angle_step_y, uint angle_offset_first_element_in_bytes)</div><div class="ttdoc">Calculate the gradient and angle from horizontal and vertical result of sobel result. </div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00121">canny.cl:121</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_a92f95668118a1bc6ef0eed965dfd570f"><div class="ttname"><a href="canny_8cl.xhtml#a92f95668118a1bc6ef0eed965dfd570f">combine_gradients_L1</a></div><div class="ttdeci">__kernel void combine_gradients_L1(__global uchar *src1_ptr, uint src1_stride_x, uint src1_step_x, uint src1_stride_y, uint src1_step_y, uint src1_offset_first_element_in_bytes, __global uchar *src2_ptr, uint src2_stride_x, uint src2_step_x, uint src2_stride_y, uint src2_step_y, uint src2_offset_first_element_in_bytes, __global uchar *grad_ptr, uint grad_stride_x, uint grad_step_x, uint grad_stride_y, uint grad_step_y, uint grad_offset_first_element_in_bytes, __global uchar *angle_ptr, uint angle_stride_x, uint angle_step_x, uint angle_stride_y, uint angle_step_y, uint angle_offset_first_element_in_bytes)</div><div class="ttdoc">Calculate the magnitude and phase from horizontal and vertical result of sobel result. </div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00057">canny.cl:57</a></div></div>
<div class="ttc" id="struct_image_xhtml_a4f0b90c9ecd6e57ceb3f37332fefe8f1"><div class="ttname"><a href="struct_image.xhtml#a4f0b90c9ecd6e57ceb3f37332fefe8f1">Image::stride_y</a></div><div class="ttdeci">int stride_y</div><div class="ttdoc">Stride of the image in Y dimension (in bytes) </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00100">helpers.h:100</a></div></div>
<div class="ttc" id="canny_8cl_xhtml_acd62ae1c9f3d7a1c7e49499d308d1904"><div class="ttname"><a href="canny_8cl.xhtml#acd62ae1c9f3d7a1c7e49499d308d1904">hysteresis</a></div><div class="ttdeci">kernel void hysteresis(__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 *out_ptr, uint out_stride_x, uint out_step_x, uint out_stride_y, uint out_step_y, uint out_offset_first_element_in_bytes, __global uchar *visited_ptr, uint visited_stride_x, uint visited_step_x, uint visited_stride_y, uint visited_step_y, uint visited_offset_first_element_in_bytes, __global uchar *recorded_ptr, uint recorded_stride_x, uint recorded_step_x, uint recorded_stride_y, uint recorded_step_y, uint recorded_offset_first_element_in_bytes, __global uchar *l1_stack_ptr, uint l1_stack_stride_x, uint l1_stack_step_x, uint l1_stack_stride_y, uint l1_stack_step_y, uint l1_stack_offset_first_element_in_bytes, __global uchar *l1_stack_counter_ptr, uint l1_stack_counter_stride_x, uint l1_stack_counter_step_x, uint l1_stack_counter_stride_y, uint l1_stack_counter_step_y, uint l1_stack_counter_offset_first_element_in_bytes, uint low_thr, uint up_thr, int width, int height)</div><div class="ttdoc">Perform hysteresis. </div><div class="ttdef"><b>Definition:</b> <a href="canny_8cl_source.xhtml#l00306">canny.cl:306</a></div></div>
<div class="ttc" id="struct_image_xhtml_ae01febbfd0689ef709f3ff6fdd2abc7e"><div class="ttname"><a href="struct_image.xhtml#ae01febbfd0689ef709f3ff6fdd2abc7e">Image::stride_x</a></div><div class="ttdeci">int stride_x</div><div class="ttdoc">Stride of the image in X dimension (in bytes) </div><div class="ttdef"><b>Definition:</b> <a href="helpers_8h_source.xhtml#l00099">helpers.h:99</a></div></div>
</div><!-- fragment --></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="canny_8cl.xhtml">canny.cl</a></li>
<li class="footer">Generated on Wed May 3 2017 17:20:04 for ARM 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>
</div>
</body>
</html>
|
