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
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
|
<!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/row_wise_reduction_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('row__wise__reduction__template_8hpp.html','');
</script>
<div id="doc-content">
<div class="header">
<div class="headertitle">
<div class="title">viennacl/device_specific/templates/row_wise_reduction_template.hpp</div> </div>
</div><!--header-->
<div class="contents">
<a href="row__wise__reduction__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_ROW_WISE_REDUCTION_HPP</span>
<a name="l00002"></a>00002 <span class="preprocessor"></span><span class="preprocessor">#define VIENNACL_DEVICE_SPECIFIC_TEMPLATES_ROW_WISE_REDUCTION_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
<a name="l00029"></a>00029 <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="l00030"></a>00030
<a name="l00031"></a>00031 <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="l00032"></a>00032 <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="l00033"></a>00033 <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="l00034"></a>00034
<a name="l00035"></a>00035 <span class="preprocessor">#include "<a class="code" href="template__base_8hpp.html">viennacl/device_specific/templates/template_base.hpp</a>"</span>
<a name="l00036"></a>00036 <span class="preprocessor">#include "<a class="code" href="device__specific_2templates_2utils_8hpp.html">viennacl/device_specific/templates/utils.hpp</a>"</span>
<a name="l00037"></a>00037
<a name="l00038"></a>00038 <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="l00039"></a>00039
<a name="l00040"></a>00040 <span class="preprocessor">#include "<a class="code" href="io_8hpp.html" title="Some helper routines for reading/writing/printing scheduler expressions.">viennacl/scheduler/io.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="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html">00047</a> <span class="keyword">struct </span><a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html">row_wise_reduction_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><a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#a64f134200cb1d4f43ad25126cfd9ddcc">00049</a> <a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#a64f134200cb1d4f43ad25126cfd9ddcc">row_wise_reduction_parameters</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _simd_width,
<a name="l00050"></a>00050 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _local_size_0, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _local_size_1,
<a name="l00051"></a>00051 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _num_groups_0, <a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7c">fetching_policy_type</a> _fetch_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, _local_size_0, _local_size_1, 1),
<a name="l00052"></a>00052 <a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#a18afc5821009c52835368d9f5ff37484">num_groups_0</a>(_num_groups_0), <a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#af86872e2bbed3201309370e46a2db318">fetch_policy</a>(_fetch_policy) { }
<a name="l00053"></a>00053
<a name="l00054"></a><a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#a18afc5821009c52835368d9f5ff37484">00054</a> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#a18afc5821009c52835368d9f5ff37484">num_groups_0</a>;
<a name="l00055"></a><a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#af86872e2bbed3201309370e46a2db318">00055</a> <a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7c">fetching_policy_type</a> <a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html#af86872e2bbed3201309370e46a2db318">fetch_policy</a>;
<a name="l00056"></a>00056 };
<a name="l00057"></a>00057
<a name="l00058"></a><a class="code" href="classviennacl_1_1device__specific_1_1row__wise__reduction__template.html">00058</a> <span class="keyword">class </span><a class="code" href="classviennacl_1_1device__specific_1_1row__wise__reduction__template.html">row_wise_reduction_template</a> : <span class="keyword">public</span> <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html">template_base_impl</a><row_wise_reduction_template, row_wise_reduction_parameters>
<a name="l00059"></a>00059 {
<a name="l00060"></a>00060 <span class="keyword">private</span>:
<a name="l00061"></a>00061 <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="l00062"></a>00062 <span class="keyword"> </span>{
<a name="l00063"></a>00063 <span class="keywordflow">if</span> (<a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.fetch_policy==<a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7ca43d469b7977b01181b327a3802b75e61">FETCH_FROM_LOCAL</a>)
<a name="l00064"></a>00064 <span class="keywordflow">return</span> TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
<a name="l00065"></a>00065 <span class="keywordflow">return</span> TEMPLATE_VALID;
<a name="l00066"></a>00066 }
<a name="l00067"></a>00067
<a name="l00068"></a>00068 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> n_lmem_elements()<span class="keyword"> const</span>
<a name="l00069"></a>00069 <span class="keyword"> </span>{
<a name="l00070"></a>00070 <span class="keywordflow">return</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>*(<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#a0013ca39734fd5a1b8110e8710a178c2">local_size_1</a>+1);
<a name="l00071"></a>00071 }
<a name="l00072"></a>00072
<a name="l00073"></a>00073 <span class="keyword">static</span> <span class="keywordtype">void</span> parse(<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, std::vector<vcl_size_t> & idx, <span class="keywordtype">bool</span> & is_trans, <a class="code" href="structviennacl_1_1scheduler_1_1lhs__rhs__element.html" title="A class representing the 'data' for the LHS or RHS operand of the respective node.">scheduler::lhs_rhs_element</a> & <a class="code" href="classviennacl_1_1matrix.html" title="A dense matrix class.">matrix</a>)
<a name="l00074"></a>00074 {
<a name="l00075"></a>00075 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a70f9b436519b02537064504f3a327470" title="Recursively execute a functor on a statement.">tree_parsing::traverse</a>(statement, statement.<a class="code" href="classviennacl_1_1scheduler_1_1statement.html#a0bf20346036ee130e78351991543cf09">root</a>(), <a class="code" href="classviennacl_1_1device__specific_1_1tree__parsing_1_1filter.html">tree_parsing::filter</a>(&<a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#acdf13a51c36290762dc8879c6541eec9">utils::is_reduction</a>, idx), <span class="keyword">false</span>);
<a name="l00076"></a>00076 is_trans = <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#aa3f7c1f119f0d1e2bc2c0e580cb8b10a">is_node_trans</a>(statement.<a class="code" href="classviennacl_1_1scheduler_1_1statement.html#afb794865f5960e34f91032674ea87cf2">array</a>(), idx[0], <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>);
<a name="l00077"></a>00077 matrix = <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>(), idx[0]).lhs;
<a name="l00078"></a>00078 }
<a name="l00079"></a>00079
<a name="l00080"></a>00080 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="keywordtype">unsigned</span> <span class="keywordtype">int</span> simd_width, <span class="keywordtype">bool</span> is_trans, std::vector<mapped_row_wise_reduction*> <span class="keyword">const</span> & exprs)<span class="keyword"> const</span>
<a name="l00081"></a>00081 <span class="keyword"> </span>{
<a name="l00082"></a>00082 <span class="keyword">using</span> <a class="code" href="namespaceviennacl_1_1scheduler_1_1detail.html#a3bd3a915c10b493d149eab2f40af2534" title="Helper routine for converting the operation enums to string.">tools::to_string</a>;
<a name="l00083"></a>00083
<a name="l00084"></a>00084 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> lsize0 = <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="l00085"></a>00085 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> lsize1 = <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#a0013ca39734fd5a1b8110e8710a178c2">local_size_1</a>+1;
<a name="l00086"></a>00086 std::string lsize1str = <a class="code" href="namespaceviennacl_1_1scheduler_1_1detail.html#a3bd3a915c10b493d149eab2f40af2534" title="Helper routine for converting the operation enums to string.">to_string</a>(lsize1);
<a name="l00087"></a>00087
<a name="l00088"></a>00088 <a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html">utils::kernel_generation_stream</a> stream;
<a name="l00089"></a>00089
<a name="l00090"></a>00090 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">","</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#a0013ca39734fd5a1b8110e8710a178c2">local_size_1</a> << <span class="stringliteral">",1)))"</span> << std::endl;
<a name="l00091"></a>00091 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a38cba6acc7e232c529d5b20b71e417f4">generate_prototype</a>(stream, kernel_prefix, <span class="stringliteral">"unsigned int M, unsigned int N,"</span>, mappings, statements);
<a name="l00092"></a>00092 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00093"></a>00093 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00094"></a>00094
<a name="l00095"></a>00095 <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="l00096"></a>00096 <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 += #start1 + #start2*#ld;"</span>, statements, mappings);
<a name="l00097"></a>00097 <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="l00098"></a>00098
<a name="l00099"></a>00099 <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">"#ld *= #nldstride;"</span>, statements, mappings);
<a name="l00100"></a>00100
<a name="l00101"></a>00101 <span class="keywordflow">for</span> (std::vector<mapped_row_wise_reduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it)
<a name="l00102"></a>00102 stream << (*it)->process(<span class="stringliteral">"__local #scalartype #name_buf["</span> + <a class="code" href="namespaceviennacl_1_1scheduler_1_1detail.html#a3bd3a915c10b493d149eab2f40af2534" title="Helper routine for converting the operation enums to string.">to_string</a>(lsize0*lsize1) + <span class="stringliteral">"];"</span>) << std::endl;
<a name="l00103"></a>00103
<a name="l00104"></a>00104 stream << <span class="stringliteral">"unsigned int lid0 = get_local_id(0);"</span> << std::endl;
<a name="l00105"></a>00105 stream << <span class="stringliteral">"unsigned int lid1 = get_local_id(1);"</span> << std::endl;
<a name="l00106"></a>00106 stream << <span class="stringliteral">"unsigned int upper_bound_0 = ( M +"</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> - 1 << <span class="stringliteral">")/"</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">"*"</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">";"</span> << std::endl;
<a name="l00107"></a>00107 stream << <span class="stringliteral">"for(unsigned int r = get_global_id(0); r < upper_bound_0; r += get_global_size(0)){"</span> << std::endl;
<a name="l00108"></a>00108 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00109"></a>00109
<a name="l00110"></a>00110 <span class="keywordflow">for</span> (std::vector<mapped_row_wise_reduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it)
<a name="l00111"></a>00111 stream << (*it)->process(<span class="stringliteral">"#scalartype #name_acc = "</span> + <a class="code" href="namespaceviennacl_1_1device__specific.html#af5ca1aac7c74718369880d50d690655e">neutral_element</a>((*it)->root_op()) + <span class="stringliteral">";"</span>) << std::endl;
<a name="l00112"></a>00112
<a name="l00113"></a>00113 stream << <span class="stringliteral">"if (r < M)"</span> << std::endl;
<a name="l00114"></a>00114 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00115"></a>00115 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00116"></a>00116
<a name="l00117"></a>00117 <span class="keyword">class </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="l00118"></a>00118 {
<a name="l00119"></a>00119 <span class="keyword">public</span>:
<a name="l00120"></a>00120 loop_body(std::vector<mapped_row_wise_reduction*> <span class="keyword">const</span> & _exprs, <span class="keywordtype">bool</span> _is_trans) : exprs(_exprs), is_trans(_is_trans){ }
<a name="l00121"></a>00121 <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> loop_simd_width)<span class="keyword"> const</span>
<a name="l00122"></a>00122 <span class="keyword"> </span>{
<a name="l00123"></a>00123 std::set<std::string> already_fetched;
<a name="l00124"></a>00124 <span class="keywordflow">for</span> (std::vector<mapped_row_wise_reduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it)
<a name="l00125"></a>00125 {
<a name="l00126"></a>00126 <span class="keywordflow">if</span> (is_trans)
<a name="l00127"></a>00127 (*it)->process_recursive(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, <span class="stringliteral">"matrix_trans"</span>, <a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#a18ef61de8cf8ec837e0107cc26f46da9">utils::append_width</a>(<span class="stringliteral">"#scalartype"</span>,loop_simd_width) + <span class="stringliteral">" #namereg = "</span> + <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a80aea1b0b07097be9834c54daf0ca765">vload</a>(loop_simd_width, <span class="stringliteral">"c*#stride1"</span>, <span class="stringliteral">"#pointer + r*#ld"</span>)+<span class="stringliteral">";"</span>, already_fetched);
<a name="l00128"></a>00128 <span class="keywordflow">else</span>
<a name="l00129"></a>00129 (*it)->process_recursive(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, <span class="stringliteral">"matrix"</span>, <span class="stringliteral">"#scalartype #namereg = #pointer[r*#stride1 + c*#ld];"</span>, already_fetched);
<a name="l00130"></a>00130 (*it)->process_recursive(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a1a3b79f66ece749edf7fc6d8586d5f1a">RHS_NODE_TYPE</a>, <span class="stringliteral">"vector"</span>, <a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#a18ef61de8cf8ec837e0107cc26f46da9">utils::append_width</a>(<span class="stringliteral">"#scalartype"</span>,loop_simd_width) + <span class="stringliteral">" #namereg = "</span> + <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a80aea1b0b07097be9834c54daf0ca765">vload</a>(loop_simd_width, <span class="stringliteral">"c*#stride"</span>, <span class="stringliteral">"#pointer"</span>)+<span class="stringliteral">";"</span>, already_fetched);
<a name="l00131"></a>00131 }
<a name="l00132"></a>00132
<a name="l00133"></a>00133
<a name="l00134"></a>00134 <span class="comment">//Update accumulators</span>
<a name="l00135"></a>00135 std::vector<std::string> str(loop_simd_width);
<a name="l00136"></a>00136 <span class="keywordflow">if</span> (loop_simd_width==1)
<a name="l00137"></a>00137 str[0] = <span class="stringliteral">"#namereg"</span>;
<a name="l00138"></a>00138 <span class="keywordflow">else</span>
<a name="l00139"></a>00139 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> a = 0; a < loop_simd_width; ++a)
<a name="l00140"></a>00140 str[a] = <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a56f2c818eebd6d26d91b0907bc7536ed">append_simd_suffix</a>(<span class="stringliteral">"#namereg.s"</span>, a);
<a name="l00141"></a>00141
<a name="l00142"></a>00142
<a name="l00143"></a>00143 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < exprs.size(); ++k)
<a name="l00144"></a>00144 {
<a name="l00145"></a>00145 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> a = 0; a < loop_simd_width; ++a)
<a name="l00146"></a>00146 {
<a name="l00147"></a>00147 std::map<std::string, std::string> accessors;
<a name="l00148"></a>00148 <span class="keywordflow">if</span> (is_trans)
<a name="l00149"></a>00149 accessors[<span class="stringliteral">"matrix_trans"</span>] = str[a];
<a name="l00150"></a>00150 <span class="keywordflow">else</span>
<a name="l00151"></a>00151 accessors[<span class="stringliteral">"matrix"</span>] = str[a];
<a name="l00152"></a>00152 accessors[<span class="stringliteral">"vector"</span>] = str[a];
<a name="l00153"></a>00153 accessors[<span class="stringliteral">"scalar"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00154"></a>00154 std::string value = exprs[k]->evaluate_recursive(<a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, accessors);
<a name="l00155"></a>00155 <span class="keywordflow">if</span> (exprs[k]->root_node().op.type==<a class="code" href="namespaceviennacl_1_1scheduler.html#a32152bde4bf97c79e26300a441099a35a3f4b97fb982d26b7d5c4dcecae48d742">scheduler::OPERATION_BINARY_MAT_VEC_PROD_TYPE</a>)
<a name="l00156"></a>00156 value+= <span class="stringliteral">"*"</span> + exprs[k]->evaluate_recursive(<a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a1a3b79f66ece749edf7fc6d8586d5f1a">RHS_NODE_TYPE</a>, accessors);
<a name="l00157"></a>00157
<a name="l00158"></a>00158 <span class="keywordflow">if</span> (exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#ae0bee04daa630b804e130cfeecbbeb70">is_index_reduction</a>())
<a name="l00159"></a>00159 <a class="code" href="namespaceviennacl_1_1device__specific.html#a9c1338a8a946a361418c94dd15809ae7">compute_index_reduction</a>(kernel_stream, exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_acc"</span>), <span class="stringliteral">"c*"</span>+<a class="code" href="namespaceviennacl_1_1scheduler_1_1detail.html#a3bd3a915c10b493d149eab2f40af2534" title="Helper routine for converting the operation enums to string.">to_string</a>(loop_simd_width) + <a class="code" href="namespaceviennacl_1_1scheduler_1_1detail.html#a3bd3a915c10b493d149eab2f40af2534" title="Helper routine for converting the operation enums to string.">to_string</a>(a), exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_acc_value"</span>), value,exprs[k]->root_op());
<a name="l00160"></a>00160 <span class="keywordflow">else</span>
<a name="l00161"></a>00161 <a class="code" href="namespaceviennacl_1_1device__specific.html#a1e2369d7f7fbcacf9c16f48cabd21b4a">compute_reduction</a>(kernel_stream, exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_acc"</span>), value,exprs[k]->root_op());
<a name="l00162"></a>00162 }
<a name="l00163"></a>00163 }
<a name="l00164"></a>00164 }
<a name="l00165"></a>00165 <span class="keyword">private</span>:
<a name="l00166"></a>00166 std::vector<mapped_row_wise_reduction*> exprs;
<a name="l00167"></a>00167 <span class="keywordtype">bool</span> is_trans;
<a name="l00168"></a>00168 };
<a name="l00169"></a>00169
<a name="l00170"></a>00170 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a72a58236e69d4882be762e72545655c1">element_wise_loop_1D</a>(stream, loop_body(exprs, is_trans), <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.fetch_policy, simd_width, <span class="stringliteral">"c"</span>, <span class="stringliteral">"N"</span>, <span class="stringliteral">"get_local_id(1)"</span>, <span class="stringliteral">"get_local_size(1)"</span>);
<a name="l00171"></a>00171 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00172"></a>00172 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00173"></a>00173
<a name="l00174"></a>00174 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < exprs.size(); ++k)
<a name="l00175"></a>00175 stream << exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_buf[lid0*"</span> + lsize1str + <span class="stringliteral">"+ lid1] = #name_acc;"</span>) << std::endl;
<a name="l00176"></a>00176
<a name="l00177"></a>00177 stream << <span class="stringliteral">"#pragma unroll"</span> << std::endl;
<a name="l00178"></a>00178 stream << <span class="stringliteral">"for(unsigned int stride = "</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#a0013ca39734fd5a1b8110e8710a178c2">local_size_1</a>/2 << <span class="stringliteral">"; stride >0; stride /=2)"</span> << std::endl;
<a name="l00179"></a>00179 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00180"></a>00180 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00181"></a>00181
<a name="l00182"></a>00182 stream << <span class="stringliteral">"barrier(CLK_LOCAL_MEM_FENCE); "</span> << std::endl;
<a name="l00183"></a>00183 stream << <span class="stringliteral">"if (lid1 < stride)"</span> << std::endl;
<a name="l00184"></a>00184 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00185"></a>00185 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00186"></a>00186
<a name="l00187"></a>00187 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < exprs.size(); k++)
<a name="l00188"></a>00188 <span class="keywordflow">if</span> (exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#ae0bee04daa630b804e130cfeecbbeb70">is_index_reduction</a>())
<a name="l00189"></a>00189 <a class="code" href="namespaceviennacl_1_1device__specific.html#a9c1338a8a946a361418c94dd15809ae7">compute_index_reduction</a>(stream, exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_buf[lid0*"</span> + lsize1str + <span class="stringliteral">" + lid1]"</span>), exprs[k]->process(<span class="stringliteral">"#name_buf[lid0*"</span> + lsize1str + <span class="stringliteral">" + lid1 + stride]"</span>)
<a name="l00190"></a>00190 , exprs[k]->process(<span class="stringliteral">"#name_buf_value[lid0*"</span> + lsize1str + <span class="stringliteral">" + lid1]"</span>), exprs[k]->process(<span class="stringliteral">"#name_buf_value[lid0*"</span> + lsize1str + <span class="stringliteral">" + lid1 + stride]"</span>),
<a name="l00191"></a>00191 exprs[k]->root_op());
<a name="l00192"></a>00192 <span class="keywordflow">else</span>
<a name="l00193"></a>00193 <a class="code" href="namespaceviennacl_1_1device__specific.html#a1e2369d7f7fbcacf9c16f48cabd21b4a">compute_reduction</a>(stream,exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_buf[lid0*"</span> + lsize1str + <span class="stringliteral">" + lid1]"</span>), exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_buf[lid0*"</span> + lsize1str + <span class="stringliteral">" + lid1 + stride]"</span>), exprs[k]->root_op());
<a name="l00194"></a>00194
<a name="l00195"></a>00195 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00196"></a>00196 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00197"></a>00197
<a name="l00198"></a>00198 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00199"></a>00199 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00200"></a>00200
<a name="l00201"></a>00201
<a name="l00202"></a>00202 stream << <span class="stringliteral">"if (lid1 == 0 && r < M)"</span>;
<a name="l00203"></a>00203 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00204"></a>00204 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00205"></a>00205 std::map<std::string, std::string> accessors;
<a name="l00206"></a>00206 accessors[<span class="stringliteral">"row_wise_reduction"</span>] = <span class="stringliteral">"#name_buf[lid0*"</span> + lsize1str + <span class="stringliteral">"]"</span>;
<a name="l00207"></a>00207 accessors[<span class="stringliteral">"vector"</span>] = <span class="stringliteral">"#pointer[r*#stride]"</span>;
<a name="l00208"></a>00208 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a810fd16fbb883a4987ff0cf720d47037">tree_parsing::evaluate</a>(stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, accessors, statements, mappings);
<a name="l00209"></a>00209 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00210"></a>00210 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00211"></a>00211
<a name="l00212"></a>00212
<a name="l00213"></a>00213 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00214"></a>00214 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00215"></a>00215
<a name="l00216"></a>00216 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00217"></a>00217 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00218"></a>00218
<a name="l00219"></a>00219 <span class="keywordflow">return</span> stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aaae555431f1866563c9ae8253bcea1db">str</a>();
<a name="l00220"></a>00220 }
<a name="l00221"></a>00221
<a name="l00222"></a>00222 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="l00223"></a>00223 <span class="keyword"> </span>{
<a name="l00224"></a>00224 std::vector<mapped_row_wise_reduction*> exprs;
<a name="l00225"></a>00225 <span class="keywordtype">bool</span> is_trans = <span class="keyword">false</span>;
<a name="l00226"></a>00226 <span class="keywordtype">bool</span> <a class="code" href="structviennacl_1_1row__major.html" title="A tag for row-major storage of a dense matrix.">row_major</a> = <span class="keyword">false</span>;
<a name="l00227"></a>00227 statements_container::data_type::const_iterator sit;
<a name="l00228"></a>00228 std::vector<mapping_type>::const_iterator mit;
<a name="l00229"></a>00229 <span class="keywordflow">for</span> (mit = mappings.begin(), sit = statements.<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html#a186e8a37443fdbc8deca9fc3f5fc64d5">data</a>().begin(); mit != mappings.end(); ++mit, ++sit)
<a name="l00230"></a>00230 {
<a name="l00231"></a>00231 std::vector<vcl_size_t> idx;
<a name="l00232"></a>00232 <a class="code" href="structviennacl_1_1scheduler_1_1lhs__rhs__element.html" title="A class representing the 'data' for the LHS or RHS operand of the respective node.">scheduler::lhs_rhs_element</a> A;
<a name="l00233"></a>00233 parse(*sit, idx, is_trans, A);
<a name="l00234"></a>00234 row_major = utils::call_on_matrix(A, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1row__major__fun.html">utils::row_major_fun</a>());
<a name="l00235"></a>00235 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> j = 0; j < idx.size(); ++j)
<a name="l00236"></a>00236 exprs.push_back((<a class="code" href="classviennacl_1_1device__specific_1_1mapped__row__wise__reduction.html" title="Vector reduction.">mapped_row_wise_reduction</a>*)(<a class="code" href="namespaceviennacl_1_1device__specific.html#ab912aba0d144f286022bb3af5d5f8f77" title="Emulation of C++11's .at() member for std::map<>, const-version.">at</a>(*mit, <a class="code" href="namespaceviennacl_1_1device__specific.html#ac45007f831a7945fb8b03ab3a8498446">mapping_key</a>(idx[j], <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>)).<span class="keyword">get</span>()));
<a name="l00237"></a>00237 }
<a name="l00238"></a>00238 is_trans = is_trans ^ <a class="code" href="namespaceviennacl_1_1traits.html#a168a8edd0b615b77f2f1da2da28f8ffb">row_major</a>;
<a name="l00239"></a>00239
<a name="l00240"></a>00240 std::vector<std::string> res;
<a name="l00241"></a>00241 <span class="keywordflow">if</span> (is_trans && <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="l00242"></a>00242 {
<a name="l00243"></a>00243 res.push_back(generate_impl(kernel_prefix, statements, mappings, <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>, is_trans, exprs));
<a name="l00244"></a>00244 res.push_back(generate_impl(kernel_prefix, statements, mappings, 1, is_trans, exprs));
<a name="l00245"></a>00245 }
<a name="l00246"></a>00246 <span class="keywordflow">else</span>
<a name="l00247"></a>00247 res.push_back(generate_impl(kernel_prefix, statements, mappings, 1, is_trans, exprs));
<a name="l00248"></a>00248
<a name="l00249"></a>00249 <span class="keywordflow">return</span> res;
<a name="l00250"></a>00250 }
<a name="l00251"></a>00251 <span class="keyword">public</span>:
<a name="l00252"></a><a class="code" href="classviennacl_1_1device__specific_1_1row__wise__reduction__template.html#ac77f0580e344b4f4f9e73a6d9c16fd5f">00252</a> <a class="code" href="classviennacl_1_1device__specific_1_1row__wise__reduction__template.html#ac77f0580e344b4f4f9e73a6d9c16fd5f">row_wise_reduction_template</a>(<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html">row_wise_reduction_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>, <span class="keywordtype">char</span> A_trans, <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_1row__wise__reduction__template.html">row_wise_reduction_template</a>, <a class="code" href="structviennacl_1_1device__specific_1_1row__wise__reduction__parameters.html">row_wise_reduction_parameters</a>>(parameters, binding_policy), A_trans_(A_trans){ }
<a name="l00253"></a>00253
<a name="l00254"></a><a class="code" href="classviennacl_1_1device__specific_1_1row__wise__reduction__template.html#a27b5b2c06f3ec88c10726a40a25c1030">00254</a> <span class="keywordtype">void</span> <a class="code" href="classviennacl_1_1device__specific_1_1row__wise__reduction__template.html#a27b5b2c06f3ec88c10726a40a25c1030">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="l00255"></a>00255 {
<a name="l00256"></a>00256 std::vector<vcl_size_t> idx;
<a name="l00257"></a>00257 <a class="code" href="structviennacl_1_1scheduler_1_1lhs__rhs__element.html" title="A class representing the 'data' for the LHS or RHS operand of the respective node.">scheduler::lhs_rhs_element</a> A;
<a name="l00258"></a>00258 <span class="keywordtype">bool</span> is_trans;
<a name="l00259"></a>00259 parse(statements.<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html#a186e8a37443fdbc8deca9fc3f5fc64d5">data</a>().front(), idx, is_trans, A);
<a name="l00260"></a>00260 <span class="keywordtype">bool</span> <a class="code" href="structviennacl_1_1row__major.html" title="A tag for row-major storage of a dense matrix.">row_major</a> = utils::call_on_matrix(A, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1row__major__fun.html">utils::row_major_fun</a>());
<a name="l00261"></a>00261
<a name="l00262"></a>00262 <a class="code" href="classviennacl_1_1ocl_1_1kernel.html" title="Represents an OpenCL kernel within ViennaCL.">viennacl::ocl::kernel</a> * kernel;
<a name="l00263"></a>00263 <span class="keywordflow">if</span> ((is_trans ^ row_major)&& <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="l00264"></a>00264 {
<a name="l00265"></a>00265 <span class="keywordflow">if</span> (<a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#ad27631755cf136243968c0548d538733">has_strided_access</a>(statements))
<a name="l00266"></a>00266 kernel = &programs[1].program().get_kernel(kernel_prefix);
<a name="l00267"></a>00267 <span class="keywordflow">else</span>
<a name="l00268"></a>00268 kernel = &programs[0].program().get_kernel(kernel_prefix);
<a name="l00269"></a>00269 }
<a name="l00270"></a>00270 <span class="keywordflow">else</span>
<a name="l00271"></a>00271 kernel = &programs[0].program().get_kernel(kernel_prefix);
<a name="l00272"></a>00272
<a name="l00273"></a>00273 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="l00274"></a>00274 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>(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#a0013ca39734fd5a1b8110e8710a178c2">local_size_1</a>);
<a name="l00275"></a>00275 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_0);
<a name="l00276"></a>00276 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>(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#a0013ca39734fd5a1b8110e8710a178c2">local_size_1</a>);
<a name="l00277"></a>00277
<a name="l00278"></a>00278 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> current_arg = 0;
<a name="l00279"></a>00279 <span class="keywordflow">if</span> (is_trans)
<a name="l00280"></a>00280 {
<a name="l00281"></a>00281 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++, cl_uint(utils::call_on_matrix(A, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1size2__fun.html">utils::size2_fun</a>())));
<a name="l00282"></a>00282 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++, cl_uint(utils::call_on_matrix(A, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1size1__fun.html">utils::size1_fun</a>())));
<a name="l00283"></a>00283 }
<a name="l00284"></a>00284 <span class="keywordflow">else</span>
<a name="l00285"></a>00285 {
<a name="l00286"></a>00286 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++, cl_uint(utils::call_on_matrix(A, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1size1__fun.html">utils::size1_fun</a>())));
<a name="l00287"></a>00287 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++, cl_uint(utils::call_on_matrix(A, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1size2__fun.html">utils::size2_fun</a>())));
<a name="l00288"></a>00288 }
<a name="l00289"></a>00289
<a name="l00290"></a>00290
<a name="l00291"></a>00291 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a313b8339e22e2a1fdd5c2c885619f1ef">set_arguments</a>(statements, *kernel, current_arg);
<a name="l00292"></a>00292 <a class="code" href="classviennacl_1_1device__specific_1_1row__wise__reduction__template.html#a27b5b2c06f3ec88c10726a40a25c1030">viennacl::ocl::enqueue</a>(*kernel);
<a name="l00293"></a>00293 }
<a name="l00294"></a>00294
<a name="l00295"></a>00295 <span class="keyword">private</span>:
<a name="l00296"></a>00296 <span class="keyword">const</span> <span class="keywordtype">char</span> A_trans_;
<a name="l00297"></a>00297 };
<a name="l00298"></a>00298
<a name="l00299"></a>00299 }
<a name="l00300"></a>00300 }
<a name="l00301"></a>00301
<a name="l00302"></a>00302 <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="row__wise__reduction__template_8hpp.html">row_wise_reduction_template.hpp</a> </li>
<li class="footer">Generated on Sat Aug 8 2015 11:59:34 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>
|
