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
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
|
<!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"/>
<title>ViennaCL - The Vienna Computing Library: viennacl/device_specific/templates/vector_axpy_template.hpp Source File</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
$(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
</head>
<body>
<div id="top"><!-- do not remove this div! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td style="padding-left: 0.5em;">
<div id="projectname">ViennaCL - The Vienna Computing Library
 <span id="projectnumber">1.7.0</span>
</div>
<div id="projectbrief">Free open-source GPU-accelerated linear algebra and solver library.</div>
</td>
<td> <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>
</td>
</tr>
</tbody>
</table>
</div>
<!-- Generated by Doxygen 1.7.6.1 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
</div>
<div id="side-nav" class="ui-resizable side-nav-resizable">
<div id="nav-tree">
<div id="nav-tree-contents">
</div>
</div>
<div id="splitbar" style="-moz-user-select:none;"
class="ui-resizable-handle">
</div>
</div>
<script type="text/javascript">
initNavTree('vector__axpy__template_8hpp.html','');
</script>
<div id="doc-content">
<div class="header">
<div class="headertitle">
<div class="title">viennacl/device_specific/templates/vector_axpy_template.hpp</div> </div>
</div><!--header-->
<div class="contents">
<a href="vector__axpy__template_8hpp.html">Go to the documentation of this file.</a><div class="fragment"><pre class="fragment"><a name="l00001"></a>00001 <span class="preprocessor">#ifndef VIENNACL_DEVICE_SPECIFIC_TEMPLATES_VECTOR_AXPY_HPP</span>
<a name="l00002"></a>00002 <span class="preprocessor"></span><span class="preprocessor">#define VIENNACL_DEVICE_SPECIFIC_TEMPLATES_VECTOR_AXPY_HPP</span>
<a name="l00003"></a>00003 <span class="preprocessor"></span>
<a name="l00004"></a>00004 <span class="comment">/* =========================================================================</span>
<a name="l00005"></a>00005 <span class="comment"> Copyright (c) 2010-2015, Institute for Microelectronics,</span>
<a name="l00006"></a>00006 <span class="comment"> Institute for Analysis and Scientific Computing,</span>
<a name="l00007"></a>00007 <span class="comment"> TU Wien.</span>
<a name="l00008"></a>00008 <span class="comment"> Portions of this software are copyright by UChicago Argonne, LLC.</span>
<a name="l00009"></a>00009 <span class="comment"></span>
<a name="l00010"></a>00010 <span class="comment"> -----------------</span>
<a name="l00011"></a>00011 <span class="comment"> ViennaCL - The Vienna Computing Library</span>
<a name="l00012"></a>00012 <span class="comment"> -----------------</span>
<a name="l00013"></a>00013 <span class="comment"></span>
<a name="l00014"></a>00014 <span class="comment"> Project Head: Karl Rupp rupp@iue.tuwien.ac.at</span>
<a name="l00015"></a>00015 <span class="comment"></span>
<a name="l00016"></a>00016 <span class="comment"> (A list of authors and contributors can be found in the manual)</span>
<a name="l00017"></a>00017 <span class="comment"></span>
<a name="l00018"></a>00018 <span class="comment"> License: MIT (X11), see file LICENSE in the base directory</span>
<a name="l00019"></a>00019 <span class="comment">============================================================================= */</span>
<a name="l00020"></a>00020
<a name="l00021"></a>00021
<a name="l00027"></a>00027 <span class="preprocessor">#include <vector></span>
<a name="l00028"></a>00028 <span class="preprocessor">#include <cmath></span>
<a name="l00029"></a>00029
<a name="l00030"></a>00030 <span class="preprocessor">#include "<a class="code" href="scheduler_2forwards_8h.html" title="Provides the datastructures for dealing with a single statement such as 'x = y + z;'.">viennacl/scheduler/forwards.h</a>"</span>
<a name="l00031"></a>00031
<a name="l00032"></a>00032 <span class="preprocessor">#include "<a class="code" href="mapped__objects_8hpp.html" title="Map ViennaCL objects to generator wrappers.">viennacl/device_specific/mapped_objects.hpp</a>"</span>
<a name="l00033"></a>00033 <span class="preprocessor">#include "<a class="code" href="tree__parsing_8hpp.html" title="Code for parsing the expression trees.">viennacl/device_specific/tree_parsing.hpp</a>"</span>
<a name="l00034"></a>00034 <span class="preprocessor">#include "<a class="code" href="device__specific_2forwards_8h.html" title="Forwards declaration.">viennacl/device_specific/forwards.h</a>"</span>
<a name="l00035"></a>00035 <span class="preprocessor">#include "<a class="code" href="device__specific_2utils_8hpp.html" title="Internal utils.">viennacl/device_specific/utils.hpp</a>"</span>
<a name="l00036"></a>00036
<a name="l00037"></a>00037 <span class="preprocessor">#include "<a class="code" href="template__base_8hpp.html">viennacl/device_specific/templates/template_base.hpp</a>"</span>
<a name="l00038"></a>00038 <span class="preprocessor">#include "<a class="code" href="device__specific_2templates_2utils_8hpp.html">viennacl/device_specific/templates/utils.hpp</a>"</span>
<a name="l00039"></a>00039
<a name="l00040"></a>00040 <span class="preprocessor">#include "<a class="code" href="tools_8hpp.html" title="Various little tools used here and there in ViennaCL.">viennacl/tools/tools.hpp</a>"</span>
<a name="l00041"></a>00041
<a name="l00042"></a>00042 <span class="keyword">namespace </span>viennacl
<a name="l00043"></a>00043 {
<a name="l00044"></a>00044 <span class="keyword">namespace </span>device_specific
<a name="l00045"></a>00045 {
<a name="l00046"></a>00046
<a name="l00047"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html">00047</a> <span class="keyword">class </span><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html">vector_axpy_parameters</a> : <span class="keyword">public</span> <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html">template_base</a>::parameters_type
<a name="l00048"></a>00048 {
<a name="l00049"></a>00049 <span class="keyword">public</span>:
<a name="l00050"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#a6b24918cb443d632108b5e963999e9d1">00050</a> <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#a6b24918cb443d632108b5e963999e9d1">vector_axpy_parameters</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _simd_width,
<a name="l00051"></a>00051 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _group_size, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _num_groups,
<a name="l00052"></a>00052 <a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7c">fetching_policy_type</a> _fetching_policy) : <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html">template_base</a>::<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html#ae4f54038c3465b5fcd440c68bd1d0701">parameters_type</a>(_simd_width, _group_size, 1, 1), <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#a2c28abc31fcc29a2af1e90144371eeee">num_groups</a>(_num_groups), <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#af40787b9ff352e215522eaba0a11279a">fetching_policy</a>(_fetching_policy){ }
<a name="l00053"></a>00053
<a name="l00054"></a>00054
<a name="l00055"></a>00055
<a name="l00056"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#a2c28abc31fcc29a2af1e90144371eeee">00056</a> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#a2c28abc31fcc29a2af1e90144371eeee">num_groups</a>;
<a name="l00057"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#af40787b9ff352e215522eaba0a11279a">00057</a> <a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7c">fetching_policy_type</a> <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html#af40787b9ff352e215522eaba0a11279a">fetching_policy</a>;
<a name="l00058"></a>00058 };
<a name="l00059"></a>00059
<a name="l00060"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html">00060</a> <span class="keyword">class </span><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html">vector_axpy_template</a> : <span class="keyword">public</span> <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html">template_base_impl</a><vector_axpy_template, vector_axpy_parameters>
<a name="l00061"></a>00061 {
<a name="l00062"></a>00062 <span class="keyword">private</span>:
<a name="l00063"></a>00063 <span class="keyword">virtual</span> <span class="keywordtype">int</span> check_invalid_impl(<a class="code" href="classviennacl_1_1ocl_1_1device.html" title="A class representing a compute device (e.g. a GPU)">viennacl::ocl::device</a> <span class="keyword">const</span> & <span class="comment">/*dev*/</span>)<span class="keyword"> const</span>
<a name="l00064"></a>00064 <span class="keyword"> </span>{
<a name="l00065"></a>00065 <span class="keywordflow">if</span> (<a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.fetching_policy==<a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7ca43d469b7977b01181b327a3802b75e61">FETCH_FROM_LOCAL</a>)
<a name="l00066"></a>00066 <span class="keywordflow">return</span> TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
<a name="l00067"></a>00067 <span class="keywordflow">return</span> TEMPLATE_VALID;
<a name="l00068"></a>00068 }
<a name="l00069"></a>00069
<a name="l00070"></a>00070 std::vector<std::string> generate_impl(std::string <span class="keyword">const</span> & kernel_prefix, <a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html">statements_container</a> <span class="keyword">const</span> & statements, std::vector<mapping_type> <span class="keyword">const</span> & mappings)<span class="keyword"> const</span>
<a name="l00071"></a>00071 <span class="keyword"> </span>{
<a name="l00072"></a>00072 std::vector<std::string> result;
<a name="l00073"></a>00073 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> i = 0; i < 2; ++i)
<a name="l00074"></a>00074 {
<a name="l00075"></a>00075 <a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html">utils::kernel_generation_stream</a> stream;
<a name="l00076"></a>00076 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> simd_width = (i==0)?1:<a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html#af859aa51f8a2c20e654f63fdb34bb5b6">simd_width</a>;
<a name="l00077"></a>00077 std::string suffix = (i==0)?<span class="stringliteral">"_strided"</span>:<span class="stringliteral">""</span>;
<a name="l00078"></a>00078 stream << <span class="stringliteral">" __attribute__((reqd_work_group_size("</span> << <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a> << <span class="stringliteral">",1,1)))"</span> << std::endl;
<a name="l00079"></a>00079 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a38cba6acc7e232c529d5b20b71e417f4">generate_prototype</a>(stream, kernel_prefix + suffix, <span class="stringliteral">"unsigned int N,"</span>, mappings, statements);
<a name="l00080"></a>00080 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00081"></a>00081 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00082"></a>00082
<a name="l00083"></a>00083 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"scalar"</span>, <span class="stringliteral">"#scalartype #namereg = *#pointer;"</span>, statements, mappings);
<a name="l00084"></a>00084 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"matrix"</span>, <span class="stringliteral">"#pointer += $OFFSET{#start1, #start2};"</span>, statements, mappings);
<a name="l00085"></a>00085 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"vector"</span>, <span class="stringliteral">"#pointer += #start;"</span>, statements, mappings);
<a name="l00086"></a>00086
<a name="l00087"></a>00087 <span class="keyword">struct </span>loop_body : <span class="keyword">public</span> <a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1loop__body__base.html">loop_body_base</a>
<a name="l00088"></a>00088 {
<a name="l00089"></a>00089 loop_body(<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html">statements_container</a> <span class="keyword">const</span> & statements_, std::vector<mapping_type> <span class="keyword">const</span> & mappings_) : statements(statements_), mappings(mappings_) { }
<a name="l00090"></a>00090
<a name="l00091"></a>00091 <span class="keywordtype">void</span> operator()(<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html">utils::kernel_generation_stream</a> & kernel_stream, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> kernel_simd_width)<span class="keyword"> const</span>
<a name="l00092"></a>00092 <span class="keyword"> </span>{
<a name="l00093"></a>00093 std::string process_str;
<a name="l00094"></a>00094 std::string i_str = (kernel_simd_width==1)?<span class="stringliteral">"i*#stride"</span>:<span class="stringliteral">"i"</span>;
<a name="l00095"></a>00095
<a name="l00096"></a>00096 process_str = <a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#a18ef61de8cf8ec837e0107cc26f46da9">utils::append_width</a>(<span class="stringliteral">"#scalartype"</span>,kernel_simd_width) + <span class="stringliteral">" #namereg = "</span> + <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a80aea1b0b07097be9834c54daf0ca765">vload</a>(kernel_simd_width, i_str, <span class="stringliteral">"#pointer"</span>) + <span class="stringliteral">";"</span>;
<a name="l00097"></a>00097 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"vector"</span>, process_str, statements, mappings);
<a name="l00098"></a>00098 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"matrix_row"</span>, <span class="stringliteral">"#scalartype #namereg = #pointer[$OFFSET{#row*#stride1, i*#stride2}];"</span>, statements, mappings);
<a name="l00099"></a>00099 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"matrix_column"</span>, <span class="stringliteral">"#scalartype #namereg = #pointer[$OFFSET{i*#stride1,#column*#stride2}];"</span>, statements, mappings);
<a name="l00100"></a>00100 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"matrix_diag"</span>, <span class="stringliteral">"#scalartype #namereg = #pointer[#diag_offset<0?$OFFSET{(i - #diag_offset)*#stride1, i*#stride2}:$OFFSET{i*#stride1, (i + #diag_offset)*#stride2}];"</span>, statements, mappings);
<a name="l00101"></a>00101
<a name="l00102"></a>00102 std::map<std::string, std::string> accessors;
<a name="l00103"></a>00103 accessors[<span class="stringliteral">"vector"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00104"></a>00104 accessors[<span class="stringliteral">"matrix_row"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00105"></a>00105 accessors[<span class="stringliteral">"matrix_column"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00106"></a>00106 accessors[<span class="stringliteral">"matrix_diag"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00107"></a>00107 accessors[<span class="stringliteral">"scalar"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00108"></a>00108 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a810fd16fbb883a4987ff0cf720d47037">tree_parsing::evaluate</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, accessors, statements, mappings);
<a name="l00109"></a>00109
<a name="l00110"></a>00110 process_str = <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#aab0c82a5d605a94281f39b69a834f1b5">vstore</a>(kernel_simd_width, <span class="stringliteral">"#namereg"</span>,i_str,<span class="stringliteral">"#pointer"</span>)+<span class="stringliteral">";"</span>;
<a name="l00111"></a>00111 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, <span class="stringliteral">"vector"</span>, process_str, statements, mappings);
<a name="l00112"></a>00112 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, <span class="stringliteral">"matrix_row"</span>, <span class="stringliteral">"#pointer[$OFFSET{#row, i}] = #namereg;"</span>, statements, mappings);
<a name="l00113"></a>00113 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, <span class="stringliteral">"matrix_column"</span>, <span class="stringliteral">"#pointer[$OFFSET{i, #column}] = #namereg;"</span>, statements, mappings);
<a name="l00114"></a>00114 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, <span class="stringliteral">"matrix_diag"</span>, <span class="stringliteral">"#pointer[#diag_offset<0?$OFFSET{i - #diag_offset, i}:$OFFSET{i, i + #diag_offset}] = #namereg;"</span>, statements, mappings);
<a name="l00115"></a>00115
<a name="l00116"></a>00116 }
<a name="l00117"></a>00117
<a name="l00118"></a>00118 <span class="keyword">private</span>:
<a name="l00119"></a>00119 <a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html">statements_container</a> <span class="keyword">const</span> & statements;
<a name="l00120"></a>00120 std::vector<mapping_type> <span class="keyword">const</span> & mappings;
<a name="l00121"></a>00121 };
<a name="l00122"></a>00122
<a name="l00123"></a>00123 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a72a58236e69d4882be762e72545655c1">element_wise_loop_1D</a>(stream, loop_body(statements, mappings), <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.fetching_policy, simd_width, <span class="stringliteral">"i"</span>, <span class="stringliteral">"N"</span>, <span class="stringliteral">"get_global_id(0)"</span>, <span class="stringliteral">"get_global_size(0)"</span>);
<a name="l00124"></a>00124
<a name="l00125"></a>00125 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00126"></a>00126 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00127"></a>00127 result.push_back(stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aaae555431f1866563c9ae8253bcea1db">str</a>());
<a name="l00128"></a>00128 }
<a name="l00129"></a>00129
<a name="l00130"></a>00130 <span class="keywordflow">return</span> result;
<a name="l00131"></a>00131 }
<a name="l00132"></a>00132
<a name="l00133"></a>00133 <span class="keyword">public</span>:
<a name="l00134"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html#a48722b720b0678f89786182b268b9e60">00134</a> <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html#a48722b720b0678f89786182b268b9e60">vector_axpy_template</a>(<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html">vector_axpy_template::parameters_type</a> <span class="keyword">const</span> & <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ad3dd3fddf4d52fd09cac0118cc993c70">parameters</a>, <a class="code" href="namespaceviennacl_1_1device__specific.html#acc7955a54ca37ad328cc400859d39be9">binding_policy_t</a> binding_policy = <a class="code" href="namespaceviennacl_1_1device__specific.html#acc7955a54ca37ad328cc400859d39be9a763b24a54e9bdb083926bc27e90ba0d6">BIND_ALL_UNIQUE</a>) : <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html">template_base_impl</a><<a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html">vector_axpy_template</a>, <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__parameters.html">vector_axpy_parameters</a>>(parameters, binding_policy), up_to_internal_size_(false){ }
<a name="l00135"></a>00135
<a name="l00136"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html#a950bc84f7f5240d3b5fe3d584c9a0bfa">00136</a> <span class="keywordtype">void</span> <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html#a950bc84f7f5240d3b5fe3d584c9a0bfa">up_to_internal_size</a>(<span class="keywordtype">bool</span> v) { up_to_internal_size_ = v; }
<a name="l00137"></a>00137
<a name="l00138"></a><a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html#a50549ae7334802a367e5568bf9e1b0ea">00138</a> <span class="keywordtype">void</span> <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html#a50549ae7334802a367e5568bf9e1b0ea">enqueue</a>(std::string <span class="keyword">const</span> & kernel_prefix, std::vector<lazy_program_compiler> & programs, <a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html">statements_container</a> <span class="keyword">const</span> & statements)
<a name="l00139"></a>00139 {
<a name="l00140"></a>00140 <a class="code" href="classviennacl_1_1ocl_1_1kernel.html" title="Represents an OpenCL kernel within ViennaCL.">viennacl::ocl::kernel</a> * kernel;
<a name="l00141"></a>00141 <span class="keywordflow">if</span> (<a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#ad27631755cf136243968c0548d538733">has_strided_access</a>(statements) && <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html#af859aa51f8a2c20e654f63fdb34bb5b6">simd_width</a> > 1)
<a name="l00142"></a>00142 kernel = &programs[0].program().get_kernel(kernel_prefix+<span class="stringliteral">"_strided"</span>);
<a name="l00143"></a>00143 <span class="keywordflow">else</span>
<a name="l00144"></a>00144 kernel = &programs[1].program().get_kernel(kernel_prefix);
<a name="l00145"></a>00145
<a name="l00146"></a>00146 kernel-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#ae2b841c487f9ddbfcfb6297d648e2d7f" title="Returns the local work size at the respective dimension.">local_work_size</a>(0, <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>);
<a name="l00147"></a>00147 kernel-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#a0144c18ae9f07722e5b5697335b7cff5" title="Returns the global work size at the respective dimension.">global_work_size</a>(0, <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>*<a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.num_groups);
<a name="l00148"></a>00148 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> current_arg = 0;
<a name="l00149"></a>00149 <a class="code" href="classviennacl_1_1scheduler_1_1statement.html" title="The main class for representing a statement such as x = inner_prod(y,z); at runtime.">scheduler::statement</a> <span class="keyword">const</span> & statement = statements.<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html#a186e8a37443fdbc8deca9fc3f5fc64d5">data</a>().front();
<a name="l00150"></a>00150 cl_uint <a class="code" href="namespaceviennacl_1_1traits.html#aa2344ea20469f55fbc15a8e9526494d0" title="Generic routine for obtaining the size of a vector (ViennaCL, uBLAS, etc.)">size</a> = <span class="keyword">static_cast<</span>cl_uint<span class="keyword">></span>(<a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a7fd9bd4940452f0e8ede97c43ba71e6a">vector_size</a>(<a class="code" href="namespaceviennacl_1_1device__specific.html#acf5956ffbc42d6804ba7f9c2376ddc75">lhs_most</a>(statement.<a class="code" href="classviennacl_1_1scheduler_1_1statement.html#afb794865f5960e34f91032674ea87cf2">array</a>(), statement.<a class="code" href="classviennacl_1_1scheduler_1_1statement.html#a0bf20346036ee130e78351991543cf09">root</a>()), up_to_internal_size_));
<a name="l00151"></a>00151 kernel-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#a86bff03b705c7cd32b1374b98477133e" title="Sets a char argument at the provided position.">arg</a>(current_arg++, size);
<a name="l00152"></a>00152 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a313b8339e22e2a1fdd5c2c885619f1ef">set_arguments</a>(statements, *kernel, current_arg);
<a name="l00153"></a>00153 <a class="code" href="classviennacl_1_1device__specific_1_1vector__axpy__template.html#a50549ae7334802a367e5568bf9e1b0ea">viennacl::ocl::enqueue</a>(*kernel);
<a name="l00154"></a>00154 }
<a name="l00155"></a>00155
<a name="l00156"></a>00156 <span class="keyword">private</span>:
<a name="l00157"></a>00157 <span class="keywordtype">bool</span> up_to_internal_size_;
<a name="l00158"></a>00158 };
<a name="l00159"></a>00159
<a name="l00160"></a>00160 }
<a name="l00161"></a>00161 }
<a name="l00162"></a>00162
<a name="l00163"></a>00163 <span class="preprocessor">#endif</span>
</pre></div></div><!-- contents -->
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark"> </span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark"> </span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark"> </span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark"> </span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark"> </span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark"> </span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark"> </span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark"> </span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark"> </span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark"> </span>Friends</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(10)"><span class="SelectionMark"> </span>Defines</a></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 id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="vector__axpy__template_8hpp.html">vector_axpy_template.hpp</a> </li>
<li class="footer">Generated on Sat Aug 8 2015 11:59:35 for ViennaCL - The Vienna Computing Library by
<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.7.6.1 </li>
</ul>
</div>
</body>
</html>
|
