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
|
<!-- 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/histogram.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.04</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('histogram_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">histogram.cl</div> </div>
</div><!--header-->
<div class="contents">
<a href="histogram_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) 2016, 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="l00026"></a><span class="lineno"><a class="line" href="histogram_8cl.xhtml#a2c8a35cfde24ca7728709200962e1a91"> 26</a></span> <span class="preprocessor">#define VATOMIC_INC16(histogram, win_pos) \</span></div><div class="line"><a name="l00027"></a><span class="lineno"> 27</span> <span class="preprocessor"> { \</span></div><div class="line"><a name="l00028"></a><span class="lineno"> 28</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s0); \</span></div><div class="line"><a name="l00029"></a><span class="lineno"> 29</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s1); \</span></div><div class="line"><a name="l00030"></a><span class="lineno"> 30</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s2); \</span></div><div class="line"><a name="l00031"></a><span class="lineno"> 31</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s3); \</span></div><div class="line"><a name="l00032"></a><span class="lineno"> 32</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s4); \</span></div><div class="line"><a name="l00033"></a><span class="lineno"> 33</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s5); \</span></div><div class="line"><a name="l00034"></a><span class="lineno"> 34</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s6); \</span></div><div class="line"><a name="l00035"></a><span class="lineno"> 35</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s7); \</span></div><div class="line"><a name="l00036"></a><span class="lineno"> 36</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s8); \</span></div><div class="line"><a name="l00037"></a><span class="lineno"> 37</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.s9); \</span></div><div class="line"><a name="l00038"></a><span class="lineno"> 38</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.sa); \</span></div><div class="line"><a name="l00039"></a><span class="lineno"> 39</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.sb); \</span></div><div class="line"><a name="l00040"></a><span class="lineno"> 40</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.sc); \</span></div><div class="line"><a name="l00041"></a><span class="lineno"> 41</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.sd); \</span></div><div class="line"><a name="l00042"></a><span class="lineno"> 42</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.se); \</span></div><div class="line"><a name="l00043"></a><span class="lineno"> 43</span> <span class="preprocessor"> atomic_inc(histogram + win_pos.sf); \</span></div><div class="line"><a name="l00044"></a><span class="lineno"> 44</span> <span class="preprocessor"> }</span></div><div class="line"><a name="l00045"></a><span class="lineno"> 45</span> </div><div class="line"><a name="l00068"></a><span class="lineno"><a class="line" href="histogram_8cl.xhtml#abc81d92c9655c4ec22fff9163b66279d"> 68</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="histogram_8cl.xhtml#abc81d92c9655c4ec22fff9163b66279d">hist_local_kernel</a>(<a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(input),</div><div class="line"><a name="l00069"></a><span class="lineno"> 69</span>  __local uint *histogram_local,</div><div class="line"><a name="l00070"></a><span class="lineno"> 70</span>  __global uint *restrict histogram,</div><div class="line"><a name="l00071"></a><span class="lineno"> 71</span>  uint num_bins,</div><div class="line"><a name="l00072"></a><span class="lineno"> 72</span>  uint <a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>,</div><div class="line"><a name="l00073"></a><span class="lineno"> 73</span>  uint range,</div><div class="line"><a name="l00074"></a><span class="lineno"> 74</span>  uint offrange)</div><div class="line"><a name="l00075"></a><span class="lineno"> 75</span> {</div><div class="line"><a name="l00076"></a><span class="lineno"> 76</span>  <a class="code" href="struct_image.xhtml">Image</a> input_buffer = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(input);</div><div class="line"><a name="l00077"></a><span class="lineno"> 77</span>  uint local_id_x = get_local_id(0);</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>  uint local_x_size = get_local_size(0);</div><div class="line"><a name="l00080"></a><span class="lineno"> 80</span> </div><div class="line"><a name="l00081"></a><span class="lineno"> 81</span>  <span class="keywordflow">if</span>(num_bins > local_x_size)</div><div class="line"><a name="l00082"></a><span class="lineno"> 82</span>  {</div><div class="line"><a name="l00083"></a><span class="lineno"> 83</span>  <span class="keywordflow">for</span>(<span class="keywordtype">int</span> i = local_id_x; i < num_bins; i += local_x_size)</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>  histogram_local[i] = 0;</div><div class="line"><a name="l00086"></a><span class="lineno"> 86</span>  }</div><div class="line"><a name="l00087"></a><span class="lineno"> 87</span>  }</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span>  <span class="keywordflow">else</span></div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span>  {</div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>  <span class="keywordflow">if</span>(local_id_x <= num_bins)</div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span>  {</div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span>  histogram_local[local_id_x] = 0;</div><div class="line"><a name="l00093"></a><span class="lineno"> 93</span>  }</div><div class="line"><a name="l00094"></a><span class="lineno"> 94</span>  }</div><div class="line"><a name="l00095"></a><span class="lineno"> 95</span> </div><div class="line"><a name="l00096"></a><span class="lineno"> 96</span>  uint16 vals = convert_uint16(vload16(0, input_buffer.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>));</div><div class="line"><a name="l00097"></a><span class="lineno"> 97</span> </div><div class="line"><a name="l00098"></a><span class="lineno"> 98</span>  uint16 win_pos = select(num_bins, ((vals - offset) * num_bins) / range, (vals >= offset && vals < offrange));</div><div class="line"><a name="l00099"></a><span class="lineno"> 99</span> </div><div class="line"><a name="l00100"></a><span class="lineno"> 100</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00101"></a><span class="lineno"> 101</span>  <a class="code" href="histogram_8cl.xhtml#a2c8a35cfde24ca7728709200962e1a91">VATOMIC_INC16</a>(histogram_local, win_pos);</div><div class="line"><a name="l00102"></a><span class="lineno"> 102</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00103"></a><span class="lineno"> 103</span> </div><div class="line"><a name="l00104"></a><span class="lineno"> 104</span>  <span class="keywordflow">if</span>(num_bins > local_x_size)</div><div class="line"><a name="l00105"></a><span class="lineno"> 105</span>  {</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span>  <span class="keywordflow">for</span>(<span class="keywordtype">int</span> i = local_id_x; i < num_bins; i += local_x_size)</div><div class="line"><a name="l00107"></a><span class="lineno"> 107</span>  {</div><div class="line"><a name="l00108"></a><span class="lineno"> 108</span>  atomic_add(histogram + i, histogram_local[i]);</div><div class="line"><a name="l00109"></a><span class="lineno"> 109</span>  }</div><div class="line"><a name="l00110"></a><span class="lineno"> 110</span>  }</div><div class="line"><a name="l00111"></a><span class="lineno"> 111</span>  <span class="keywordflow">else</span></div><div class="line"><a name="l00112"></a><span class="lineno"> 112</span>  {</div><div class="line"><a name="l00113"></a><span class="lineno"> 113</span>  <span class="keywordflow">if</span>(local_id_x <= num_bins)</div><div class="line"><a name="l00114"></a><span class="lineno"> 114</span>  {</div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span>  atomic_add(histogram + local_id_x, histogram_local[local_id_x]);</div><div class="line"><a name="l00116"></a><span class="lineno"> 116</span>  }</div><div class="line"><a name="l00117"></a><span class="lineno"> 117</span>  }</div><div class="line"><a name="l00118"></a><span class="lineno"> 118</span> }</div><div class="line"><a name="l00119"></a><span class="lineno"> 119</span> </div><div class="line"><a name="l00141"></a><span class="lineno"><a class="line" href="histogram_8cl.xhtml#af82fea967051b827585009463255262d"> 141</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="histogram_8cl.xhtml#af82fea967051b827585009463255262d">hist_border_kernel</a>(<a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(input),</div><div class="line"><a name="l00142"></a><span class="lineno"> 142</span>  __global uint *restrict histogram,</div><div class="line"><a name="l00143"></a><span class="lineno"> 143</span>  uint num_bins,</div><div class="line"><a name="l00144"></a><span class="lineno"> 144</span>  uint <a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>,</div><div class="line"><a name="l00145"></a><span class="lineno"> 145</span>  uint range,</div><div class="line"><a name="l00146"></a><span class="lineno"> 146</span>  uint offrange)</div><div class="line"><a name="l00147"></a><span class="lineno"> 147</span> {</div><div class="line"><a name="l00148"></a><span class="lineno"> 148</span>  <a class="code" href="struct_image.xhtml">Image</a> input_buffer = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(input);</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>  uint val = (uint)(*input_buffer.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00151"></a><span class="lineno"> 151</span> </div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span>  uint win_pos = (val >= <a class="code" href="helpers_8h.xhtml#a009469e4d9b8fce3b6d5e97d2077827d">offset</a>) ? (((val - offset) * num_bins) / range) : 0;</div><div class="line"><a name="l00153"></a><span class="lineno"> 153</span> </div><div class="line"><a name="l00154"></a><span class="lineno"> 154</span>  <span class="keywordflow">if</span>(val >= offset && (val < offrange))</div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span>  {</div><div class="line"><a name="l00156"></a><span class="lineno"> 156</span>  atomic_inc(histogram + win_pos);</div><div class="line"><a name="l00157"></a><span class="lineno"> 157</span>  }</div><div class="line"><a name="l00158"></a><span class="lineno"> 158</span> }</div><div class="line"><a name="l00159"></a><span class="lineno"> 159</span> </div><div class="line"><a name="l00178"></a><span class="lineno"><a class="line" href="histogram_8cl.xhtml#a7c8051ab952a597e66090d77f4dc60e4"> 178</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="histogram_8cl.xhtml#a7c8051ab952a597e66090d77f4dc60e4">hist_local_kernel_fixed</a>(<a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(input),</div><div class="line"><a name="l00179"></a><span class="lineno"> 179</span>  __local uint *histogram_local,</div><div class="line"><a name="l00180"></a><span class="lineno"> 180</span>  __global uint *restrict histogram)</div><div class="line"><a name="l00181"></a><span class="lineno"> 181</span> {</div><div class="line"><a name="l00182"></a><span class="lineno"> 182</span>  <a class="code" href="struct_image.xhtml">Image</a> input_buffer = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(input);</div><div class="line"><a name="l00183"></a><span class="lineno"> 183</span> </div><div class="line"><a name="l00184"></a><span class="lineno"> 184</span>  uint local_index = get_local_id(0);</div><div class="line"><a name="l00185"></a><span class="lineno"> 185</span>  uint local_x_size = get_local_size(0);</div><div class="line"><a name="l00186"></a><span class="lineno"> 186</span> </div><div class="line"><a name="l00187"></a><span class="lineno"> 187</span>  <span class="keywordflow">for</span>(<span class="keywordtype">int</span> i = local_index; i < 256; i += local_x_size)</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>  {</div><div class="line"><a name="l00189"></a><span class="lineno"> 189</span>  histogram_local[i] = 0;</div><div class="line"><a name="l00190"></a><span class="lineno"> 190</span>  }</div><div class="line"><a name="l00191"></a><span class="lineno"> 191</span> </div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>  uint16 vals = convert_uint16(vload16(0, input_buffer.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>));</div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span> </div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>  barrier(CLK_LOCAL_MEM_FENCE);</div><div class="line"><a name="l00195"></a><span class="lineno"> 195</span> </div><div class="line"><a name="l00196"></a><span class="lineno"> 196</span>  atomic_inc(histogram_local + vals.s0);</div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>  atomic_inc(histogram_local + vals.s1);</div><div class="line"><a name="l00198"></a><span class="lineno"> 198</span>  atomic_inc(histogram_local + vals.s2);</div><div class="line"><a name="l00199"></a><span class="lineno"> 199</span>  atomic_inc(histogram_local + vals.s3);</div><div class="line"><a name="l00200"></a><span class="lineno"> 200</span>  atomic_inc(histogram_local + vals.s4);</div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span>  atomic_inc(histogram_local + vals.s5);</div><div class="line"><a name="l00202"></a><span class="lineno"> 202</span>  atomic_inc(histogram_local + vals.s6);</div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>  atomic_inc(histogram_local + vals.s7);</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</span>  atomic_inc(histogram_local + vals.s8);</div><div class="line"><a name="l00205"></a><span class="lineno"> 205</span>  atomic_inc(histogram_local + vals.s9);</div><div class="line"><a name="l00206"></a><span class="lineno"> 206</span>  atomic_inc(histogram_local + vals.sa);</div><div class="line"><a name="l00207"></a><span class="lineno"> 207</span>  atomic_inc(histogram_local + vals.sb);</div><div class="line"><a name="l00208"></a><span class="lineno"> 208</span>  atomic_inc(histogram_local + vals.sc);</div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span>  atomic_inc(histogram_local + vals.sd);</div><div class="line"><a name="l00210"></a><span class="lineno"> 210</span>  atomic_inc(histogram_local + vals.se);</div><div class="line"><a name="l00211"></a><span class="lineno"> 211</span>  atomic_inc(histogram_local + vals.sf);</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span> </div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>  barrier(CLK_LOCAL_MEM_FENCE);</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="keywordflow">for</span>(<span class="keywordtype">int</span> i = local_index; i < 256; i += local_x_size)</div><div class="line"><a name="l00216"></a><span class="lineno"> 216</span>  {</div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>  atomic_add(histogram + i, histogram_local[i]);</div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>  }</div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span> }</div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span> </div><div class="line"><a name="l00238"></a><span class="lineno"><a class="line" href="histogram_8cl.xhtml#aec6ec6157573195df9694109ebbb38ae"> 238</a></span> __kernel <span class="keywordtype">void</span> <a class="code" href="histogram_8cl.xhtml#aec6ec6157573195df9694109ebbb38ae">hist_border_kernel_fixed</a>(<a class="code" href="helpers_8h.xhtml#a22f42fcf2077d951271df83b55c1a71a">IMAGE_DECLARATION</a>(input),</div><div class="line"><a name="l00239"></a><span class="lineno"> 239</span>  __global uint *restrict histogram)</div><div class="line"><a name="l00240"></a><span class="lineno"> 240</span> {</div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>  <a class="code" href="struct_image.xhtml">Image</a> input_buffer = <a class="code" href="helpers_8h.xhtml#aebe814363556c244be043b13e7969197">CONVERT_TO_IMAGE_STRUCT</a>(input);</div><div class="line"><a name="l00242"></a><span class="lineno"> 242</span>  atomic_inc(histogram + *input_buffer.<a class="code" href="struct_image.xhtml#acf52c23cbd7424606c10a606524e3e32">ptr</a>);</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span> }</div><div class="ttc" id="histogram_8cl_xhtml_a2c8a35cfde24ca7728709200962e1a91"><div class="ttname"><a href="histogram_8cl.xhtml#a2c8a35cfde24ca7728709200962e1a91">VATOMIC_INC16</a></div><div class="ttdeci">#define VATOMIC_INC16(histogram, win_pos)</div><div class="ttdef"><b>Definition:</b> <a href="histogram_8cl_source.xhtml#l00026">histogram.cl:26</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="histogram_8cl_xhtml_abc81d92c9655c4ec22fff9163b66279d"><div class="ttname"><a href="histogram_8cl.xhtml#abc81d92c9655c4ec22fff9163b66279d">hist_local_kernel</a></div><div class="ttdeci">__kernel void hist_local_kernel(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __local uint *histogram_local, __global uint *restrict histogram, uint num_bins, uint offset, uint range, uint offrange)</div><div class="ttdoc">Calculate the histogram of an 8 bit grayscale image. </div><div class="ttdef"><b>Definition:</b> <a href="histogram_8cl_source.xhtml#l00068">histogram.cl:68</a></div></div>
<div class="ttc" id="histogram_8cl_xhtml_a7c8051ab952a597e66090d77f4dc60e4"><div class="ttname"><a href="histogram_8cl.xhtml#a7c8051ab952a597e66090d77f4dc60e4">hist_local_kernel_fixed</a></div><div class="ttdeci">__kernel void hist_local_kernel_fixed(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __local uint *histogram_local, __global uint *restrict histogram)</div><div class="ttdoc">Calculate the histogram of an 8 bit grayscale image with bin size of 256 and window size of 1...</div><div class="ttdef"><b>Definition:</b> <a href="histogram_8cl_source.xhtml#l00178">histogram.cl:178</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="histogram_8cl_xhtml_af82fea967051b827585009463255262d"><div class="ttname"><a href="histogram_8cl.xhtml#af82fea967051b827585009463255262d">hist_border_kernel</a></div><div class="ttdeci">__kernel void hist_border_kernel(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __global uint *restrict histogram, uint num_bins, uint offset, uint range, uint offrange)</div><div class="ttdoc">Calculate the histogram of an 8 bit grayscale image&#39;s border. </div><div class="ttdef"><b>Definition:</b> <a href="histogram_8cl_source.xhtml#l00141">histogram.cl:141</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="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="histogram_8cl_xhtml_aec6ec6157573195df9694109ebbb38ae"><div class="ttname"><a href="histogram_8cl.xhtml#aec6ec6157573195df9694109ebbb38ae">hist_border_kernel_fixed</a></div><div class="ttdeci">__kernel void hist_border_kernel_fixed(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __global uint *restrict histogram)</div><div class="ttdoc">Calculate the histogram of an 8 bit grayscale image with bin size as 256 and window size as 1...</div><div class="ttdef"><b>Definition:</b> <a href="histogram_8cl_source.xhtml#l00238">histogram.cl:238</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="histogram_8cl.xhtml">histogram.cl</a></li>
<li class="footer">Generated on Wed Apr 12 2017 14:26:05 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>
|
