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
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
|
<!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/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('reduction__template_8hpp.html','');
</script>
<div id="doc-content">
<div class="header">
<div class="headertitle">
<div class="title">viennacl/device_specific/templates/reduction_template.hpp</div> </div>
</div><!--header-->
<div class="contents">
<a href="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_REDUCTION_TEMPLATE_HPP</span>
<a name="l00002"></a>00002 <span class="preprocessor"></span><span class="preprocessor">#define VIENNACL_DEVICE_SPECIFIC_TEMPLATES_REDUCTION_TEMPLATE_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="opencl_8hpp.html" title="Implementations for the OpenCL backend functionality.">viennacl/backend/opencl.hpp</a>"</span>
<a name="l00030"></a>00030
<a name="l00031"></a>00031 <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="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="keyword">namespace </span>viennacl
<a name="l00041"></a>00041 {
<a name="l00042"></a>00042 <span class="keyword">namespace </span>device_specific
<a name="l00043"></a>00043 {
<a name="l00044"></a>00044
<a name="l00045"></a><a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html">00045</a> <span class="keyword">struct </span><a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html">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="l00046"></a>00046 {
<a name="l00047"></a><a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#abf01b27ceb1be3066c5436d08c51a34e">00047</a> <a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#abf01b27ceb1be3066c5436d08c51a34e">reduction_parameters</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _simd_width,
<a name="l00048"></a>00048 <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="l00049"></a>00049 <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, 2), <a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#a539b1a9c7d202a007c70b4bc03ba1232">num_groups</a>(_num_groups), <a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#a9e1dd78718fc464c39afc7b2c3547179">fetching_policy</a>(_fetching_policy){ }
<a name="l00050"></a>00050
<a name="l00051"></a><a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#a539b1a9c7d202a007c70b4bc03ba1232">00051</a> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#a539b1a9c7d202a007c70b4bc03ba1232">num_groups</a>;
<a name="l00052"></a><a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#a9e1dd78718fc464c39afc7b2c3547179">00052</a> <a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7c">fetching_policy_type</a> <a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html#a9e1dd78718fc464c39afc7b2c3547179">fetching_policy</a>;
<a name="l00053"></a>00053 };
<a name="l00054"></a>00054
<a name="l00055"></a><a class="code" href="classviennacl_1_1device__specific_1_1reduction__template.html">00055</a> <span class="keyword">class </span><a class="code" href="classviennacl_1_1device__specific_1_1reduction__template.html">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><reduction_template, reduction_parameters>
<a name="l00056"></a>00056 {
<a name="l00057"></a>00057
<a name="l00058"></a>00058 <span class="keyword">private</span>:
<a name="l00059"></a>00059 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> n_lmem_elements()<span class="keyword"> const</span>
<a name="l00060"></a>00060 <span class="keyword"> </span>{
<a name="l00061"></a>00061 <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 name="l00062"></a>00062 }
<a name="l00063"></a>00063
<a name="l00064"></a>00064 <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="l00065"></a>00065 <span class="keyword"> </span>{
<a name="l00066"></a>00066 <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="l00067"></a>00067 <span class="keywordflow">return</span> TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
<a name="l00068"></a>00068 <span class="keywordflow">return</span> TEMPLATE_VALID;
<a name="l00069"></a>00069 }
<a name="l00070"></a>00070
<a name="l00071"></a>00071 <span class="keyword">inline</span> <span class="keywordtype">void</span> reduce_1d_local_memory(<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html">utils::kernel_generation_stream</a> & stream, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="namespaceviennacl_1_1traits.html#aa2344ea20469f55fbc15a8e9526494d0" title="Generic routine for obtaining the size of a vector (ViennaCL, uBLAS, etc.)">size</a>, std::vector<mapped_scalar_reduction*> exprs,
<a name="l00072"></a>00072 std::string <span class="keyword">const</span> & buf_str, std::string <span class="keyword">const</span> & buf_value_str)<span class="keyword"> const</span>
<a name="l00073"></a>00073 <span class="keyword"> </span>{
<a name="l00074"></a>00074 stream << <span class="stringliteral">"#pragma unroll"</span> << std::endl;
<a name="l00075"></a>00075 stream << <span class="stringliteral">"for(unsigned int stride = "</span> << size/2 << <span class="stringliteral">"; stride >0; stride /=2)"</span> << std::endl;
<a name="l00076"></a>00076 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00077"></a>00077 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00078"></a>00078 stream << <span class="stringliteral">"barrier(CLK_LOCAL_MEM_FENCE); "</span> << std::endl;
<a name="l00079"></a>00079 stream << <span class="stringliteral">"if (lid < stride)"</span> << std::endl;
<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 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < exprs.size(); k++)
<a name="l00084"></a>00084 <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="l00085"></a>00085 <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>(buf_str+<span class="stringliteral">"[lid]"</span>), exprs[k]->process(buf_str+<span class="stringliteral">"[lid+stride]"</span>)
<a name="l00086"></a>00086 , exprs[k]->process(buf_value_str+<span class="stringliteral">"[lid]"</span>), exprs[k]->process(buf_value_str+<span class="stringliteral">"[lid+stride]"</span>),
<a name="l00087"></a>00087 exprs[k]->root_op());
<a name="l00088"></a>00088 <span class="keywordflow">else</span>
<a name="l00089"></a>00089 <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>(buf_str+<span class="stringliteral">"[lid]"</span>), exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(buf_str+<span class="stringliteral">"[lid+stride]"</span>), exprs[k]->root_op());
<a name="l00090"></a>00090 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00091"></a>00091 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00092"></a>00092 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00093"></a>00093 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00094"></a>00094 }
<a name="l00095"></a>00095
<a name="l00096"></a>00096 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="keyword"> const</span>
<a name="l00097"></a>00097 <span class="keyword"> </span>{
<a name="l00098"></a>00098 <a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html">utils::kernel_generation_stream</a> stream;
<a name="l00099"></a>00099
<a name="l00100"></a>00100 std::vector<mapped_scalar_reduction*> exprs;
<a name="l00101"></a>00101 <span class="keywordflow">for</span> (std::vector<mapping_type>::const_iterator it = mappings.begin(); it != mappings.end(); ++it)
<a name="l00102"></a>00102 <span class="keywordflow">for</span> (mapping_type::const_iterator iit = it->begin(); iit != it->end(); ++iit)
<a name="l00103"></a>00103 <span class="keywordflow">if</span> (<a class="code" href="classviennacl_1_1device__specific_1_1mapped__scalar__reduction.html" title="Scalar reduction.">mapped_scalar_reduction</a> * p = dynamic_cast<mapped_scalar_reduction*>(iit->second.get()))
<a name="l00104"></a>00104 exprs.push_back(p);
<a name="l00105"></a>00105 <a class="code" href="namespaceviennacl.html#a98a0afcc513111ffa0bd138f891930df">vcl_size_t</a> N = exprs.size();
<a name="l00106"></a>00106
<a name="l00107"></a>00107 std::string arguments = generate_value_kernel_argument(<span class="stringliteral">"unsigned int"</span>, <span class="stringliteral">"N"</span>);
<a name="l00108"></a>00108 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < N; ++k)
<a name="l00109"></a>00109 {
<a name="l00110"></a>00110 std::string numeric_type = utils::numeric_type_to_string(<a class="code" href="namespaceviennacl_1_1device__specific.html#acf5956ffbc42d6804ba7f9c2376ddc75">lhs_most</a>(exprs[k]->statement().array(),
<a name="l00111"></a>00111 exprs[k]->statement().root()).lhs.numeric_type);
<a name="l00112"></a>00112 <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="l00113"></a>00113 {
<a name="l00114"></a>00114 arguments += generate_pointer_kernel_argument(<span class="stringliteral">"__global"</span>, <span class="stringliteral">"unsigned int"</span>, exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_temp"</span>));
<a name="l00115"></a>00115 arguments += generate_pointer_kernel_argument(<span class="stringliteral">"__global"</span>, numeric_type, exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_temp_value"</span>));
<a name="l00116"></a>00116 }
<a name="l00117"></a>00117 <span class="keywordflow">else</span>
<a name="l00118"></a>00118 arguments += generate_pointer_kernel_argument(<span class="stringliteral">"__global"</span>, numeric_type, exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_temp"</span>));
<a name="l00119"></a>00119 }
<a name="l00120"></a>00120
<a name="l00121"></a>00121
<a name="l00122"></a>00122 <span class="comment">/* ------------------------</span>
<a name="l00123"></a>00123 <span class="comment"> * First Kernel</span>
<a name="l00124"></a>00124 <span class="comment"> * -----------------------*/</span>
<a name="l00125"></a>00125 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="l00126"></a>00126 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a38cba6acc7e232c529d5b20b71e417f4">generate_prototype</a>(stream, kernel_prefix + <span class="stringliteral">"_0"</span>, arguments, mappings, statements);
<a name="l00127"></a>00127 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00128"></a>00128 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00129"></a>00129
<a name="l00130"></a>00130 <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="l00131"></a>00131 stream << <span class="stringliteral">"unsigned int lid = get_local_id(0);"</span> << std::endl;
<a name="l00132"></a>00132 <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="l00133"></a>00133
<a name="l00134"></a>00134 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < N; ++k)
<a name="l00135"></a>00135 {
<a name="l00136"></a>00136 <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="l00137"></a>00137 {
<a name="l00138"></a>00138 stream << exprs[k]->process(<span class="stringliteral">"__local #scalartype #name_buf_value["</span> + <a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::to_string</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#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>) + <span class="stringliteral">"];"</span>) << std::endl;
<a name="l00139"></a>00139 stream << exprs[k]->process(<span class="stringliteral">"#scalartype #name_acc_value = "</span> + <a class="code" href="namespaceviennacl_1_1device__specific.html#af5ca1aac7c74718369880d50d690655e">neutral_element</a>(exprs[k]->root_op()) + <span class="stringliteral">";"</span>) << std::endl;
<a name="l00140"></a>00140 stream << exprs[k]->process(<span class="stringliteral">"__local unsigned int #name_buf["</span> + <a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::to_string</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#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>) + <span class="stringliteral">"];"</span>) << std::endl;
<a name="l00141"></a>00141 stream << exprs[k]->process(<span class="stringliteral">"unsigned int #name_acc = 0;"</span>) << std::endl;
<a name="l00142"></a>00142 }
<a name="l00143"></a>00143 <span class="keywordflow">else</span>
<a name="l00144"></a>00144 {
<a name="l00145"></a>00145 stream << exprs[k]->process(<span class="stringliteral">"__local #scalartype #name_buf["</span> + <a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::to_string</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#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>) + <span class="stringliteral">"];"</span>) << std::endl;
<a name="l00146"></a>00146 stream << exprs[k]->process(<span class="stringliteral">"#scalartype #name_acc = "</span> + <a class="code" href="namespaceviennacl_1_1device__specific.html#af5ca1aac7c74718369880d50d690655e">neutral_element</a>(exprs[k]->root_op()) + <span class="stringliteral">";"</span>) << std::endl;
<a name="l00147"></a>00147 }
<a name="l00148"></a>00148 }
<a name="l00149"></a>00149
<a name="l00150"></a>00150 <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="l00151"></a>00151 {
<a name="l00152"></a>00152 <span class="keyword">public</span>:
<a name="l00153"></a>00153 loop_body(std::vector<mapped_scalar_reduction*> <span class="keyword">const</span> & exprs_) : exprs(exprs_){ }
<a name="l00154"></a>00154
<a name="l00155"></a>00155 <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="l00156"></a>00156 <span class="keyword"> </span>{
<a name="l00157"></a>00157 std::string i = (loop_simd_width==1)?<span class="stringliteral">"i*#stride"</span>:<span class="stringliteral">"i"</span>;
<a name="l00158"></a>00158 std::string process_str;
<a name="l00159"></a>00159 <span class="comment">//Fetch vector entry</span>
<a name="l00160"></a>00160 {
<a name="l00161"></a>00161 std::set<std::string> already_fetched;
<a name="l00162"></a>00162 process_str = <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,i,<span class="stringliteral">"#pointer"</span>)+<span class="stringliteral">";"</span>;
<a name="l00163"></a>00163 <span class="keywordflow">for</span> (std::vector<mapped_scalar_reduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it)
<a name="l00164"></a>00164 {
<a name="l00165"></a>00165 (*it)->process_recursive(kernel_stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"vector"</span>, process_str, already_fetched);
<a name="l00166"></a>00166 (*it)->process_recursive(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>, already_fetched);
<a name="l00167"></a>00167 (*it)->process_recursive(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>, already_fetched);
<a name="l00168"></a>00168 (*it)->process_recursive(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>, already_fetched);
<a name="l00169"></a>00169 }
<a name="l00170"></a>00170 }
<a name="l00171"></a>00171
<a name="l00172"></a>00172 <span class="comment">//Update accumulators</span>
<a name="l00173"></a>00173 std::vector<std::string> str(loop_simd_width);
<a name="l00174"></a>00174 <span class="keywordflow">if</span> (loop_simd_width==1)
<a name="l00175"></a>00175 str[0] = <span class="stringliteral">"#namereg"</span>;
<a name="l00176"></a>00176 <span class="keywordflow">else</span>
<a name="l00177"></a>00177 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> a = 0; a < loop_simd_width; ++a)
<a name="l00178"></a>00178 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="l00179"></a>00179
<a name="l00180"></a>00180 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < exprs.size(); ++k)
<a name="l00181"></a>00181 {
<a name="l00182"></a>00182 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> a = 0; a < loop_simd_width; ++a)
<a name="l00183"></a>00183 {
<a name="l00184"></a>00184 std::map<std::string, std::string> accessors;
<a name="l00185"></a>00185 accessors[<span class="stringliteral">"vector"</span>] = str[a];
<a name="l00186"></a>00186 accessors[<span class="stringliteral">"matrix_row"</span>] = str[a];
<a name="l00187"></a>00187 accessors[<span class="stringliteral">"matrix_column"</span>] = str[a];
<a name="l00188"></a>00188 accessors[<span class="stringliteral">"matrix_diag"</span>] = str[a];
<a name="l00189"></a>00189 accessors[<span class="stringliteral">"scalar"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00190"></a>00190 std::string value = exprs[k]->evaluate_recursive(<a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, accessors);
<a name="l00191"></a>00191 <span class="keywordflow">if</span> (exprs[k]->root_node().op.type==<a class="code" href="namespaceviennacl_1_1scheduler.html#a32152bde4bf97c79e26300a441099a35acf361c734fac5acbaa6ff89b838702f3">scheduler::OPERATION_BINARY_INNER_PROD_TYPE</a>)
<a name="l00192"></a>00192 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="l00193"></a>00193
<a name="l00194"></a>00194 <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="l00195"></a>00195 <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">"i*"</span>+<a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::to_string</a>(loop_simd_width) + <span class="stringliteral">"+"</span> + <a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::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="l00196"></a>00196 <span class="keywordflow">else</span>
<a name="l00197"></a>00197 <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="l00198"></a>00198 }
<a name="l00199"></a>00199 }
<a name="l00200"></a>00200 }
<a name="l00201"></a>00201
<a name="l00202"></a>00202 <span class="keyword">private</span>:
<a name="l00203"></a>00203 std::vector<mapped_scalar_reduction*> exprs;
<a name="l00204"></a>00204 };
<a name="l00205"></a>00205
<a name="l00206"></a>00206 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a72a58236e69d4882be762e72545655c1">element_wise_loop_1D</a>(stream, loop_body(exprs), <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="l00207"></a>00207
<a name="l00208"></a>00208 <span class="comment">//Fills local memory</span>
<a name="l00209"></a>00209 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < N; ++k)
<a name="l00210"></a>00210 {
<a name="l00211"></a>00211 <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="l00212"></a>00212 stream << exprs[k]->process(<span class="stringliteral">"#name_buf_value[lid] = #name_acc_value;"</span>) << std::endl;
<a name="l00213"></a>00213 stream << exprs[k]->process(<span class="stringliteral">"#name_buf[lid] = #name_acc;"</span>) << std::endl;
<a name="l00214"></a>00214 }
<a name="l00215"></a>00215
<a name="l00216"></a>00216 <span class="comment">//Reduce local memory</span>
<a name="l00217"></a>00217 reduce_1d_local_memory(stream, <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>, exprs, <span class="stringliteral">"#name_buf"</span>, <span class="stringliteral">"#name_buf_value"</span>);
<a name="l00218"></a>00218
<a name="l00219"></a>00219 <span class="comment">//Write to temporary buffers</span>
<a name="l00220"></a>00220 stream << <span class="stringliteral">"if (lid==0)"</span> << std::endl;
<a name="l00221"></a>00221 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00222"></a>00222 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00223"></a>00223 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < N; ++k)
<a name="l00224"></a>00224 {
<a name="l00225"></a>00225 <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="l00226"></a>00226 stream << exprs[k]->process(<span class="stringliteral">"#name_temp_value[get_group_id(0)] = #name_buf_value[0];"</span>) << std::endl;
<a name="l00227"></a>00227 stream << exprs[k]->process(<span class="stringliteral">"#name_temp[get_group_id(0)] = #name_buf[0];"</span>) << std::endl;
<a name="l00228"></a>00228 }
<a name="l00229"></a>00229 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00230"></a>00230 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00231"></a>00231
<a name="l00232"></a>00232 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00233"></a>00233 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00234"></a>00234
<a name="l00235"></a>00235 <span class="comment">/* ------------------------</span>
<a name="l00236"></a>00236 <span class="comment"> * Second kernel</span>
<a name="l00237"></a>00237 <span class="comment"> * -----------------------*/</span>
<a name="l00238"></a>00238 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="l00239"></a>00239 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a38cba6acc7e232c529d5b20b71e417f4">generate_prototype</a>(stream, kernel_prefix + <span class="stringliteral">"_1"</span>, arguments, mappings, statements);
<a name="l00240"></a>00240 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00241"></a>00241 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00242"></a>00242
<a name="l00243"></a>00243 stream << <span class="stringliteral">"unsigned int lid = get_local_id(0);"</span> << std::endl;
<a name="l00244"></a>00244
<a name="l00245"></a>00245 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < N; ++k)
<a name="l00246"></a>00246 {
<a name="l00247"></a>00247 <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="l00248"></a>00248 {
<a name="l00249"></a>00249 stream << exprs[k]->process(<span class="stringliteral">"__local unsigned int #name_buf["</span> + <a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::to_string</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#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>) + <span class="stringliteral">"];"</span>);
<a name="l00250"></a>00250 stream << exprs[k]->process(<span class="stringliteral">"unsigned int #name_acc = 0;"</span>) << std::endl;
<a name="l00251"></a>00251 stream << exprs[k]->process(<span class="stringliteral">"__local #scalartype #name_buf_value["</span> + <a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::to_string</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#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>) + <span class="stringliteral">"];"</span>) << std::endl;
<a name="l00252"></a>00252 stream << exprs[k]->process(<span class="stringliteral">"#scalartype #name_acc_value = "</span> + <a class="code" href="namespaceviennacl_1_1device__specific.html#af5ca1aac7c74718369880d50d690655e">neutral_element</a>(exprs[k]->root_op()) + <span class="stringliteral">";"</span>);
<a name="l00253"></a>00253 }
<a name="l00254"></a>00254 <span class="keywordflow">else</span>
<a name="l00255"></a>00255 {
<a name="l00256"></a>00256 stream << exprs[k]->process(<span class="stringliteral">"__local #scalartype #name_buf["</span> + <a class="code" href="namespaceviennacl_1_1tools.html#a534ecc78073a72515aff64c39d27b3a6">tools::to_string</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#ad8356d07e99690dafe3d12041cbcd169">local_size_0</a>) + <span class="stringliteral">"];"</span>) << std::endl;
<a name="l00257"></a>00257 stream << exprs[k]->process(<span class="stringliteral">"#scalartype #name_acc = "</span> + <a class="code" href="namespaceviennacl_1_1device__specific.html#af5ca1aac7c74718369880d50d690655e">neutral_element</a>(exprs[k]->root_op()) + <span class="stringliteral">";"</span>);
<a name="l00258"></a>00258 }
<a name="l00259"></a>00259 }
<a name="l00260"></a>00260
<a name="l00261"></a>00261 stream << <span class="stringliteral">"for(unsigned int i = lid; i < "</span> << <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.num_groups << <span class="stringliteral">"; i += get_local_size(0))"</span> << std::endl;
<a name="l00262"></a>00262 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00263"></a>00263 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00264"></a>00264 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < N; ++k)
<a name="l00265"></a>00265 <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="l00266"></a>00266 <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_acc"</span>), exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_temp[i]"</span>),
<a name="l00267"></a>00267 exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_acc_value"</span>),exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_temp_value[i]"</span>),exprs[k]->root_op());
<a name="l00268"></a>00268 <span class="keywordflow">else</span>
<a name="l00269"></a>00269 <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_acc"</span>), exprs[k]-><a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">process</a>(<span class="stringliteral">"#name_temp[i]"</span>), exprs[k]->root_op());
<a name="l00270"></a>00270
<a name="l00271"></a>00271 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00272"></a>00272 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00273"></a>00273
<a name="l00274"></a>00274 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < N; ++k)
<a name="l00275"></a>00275 {
<a name="l00276"></a>00276 <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="l00277"></a>00277 stream << exprs[k]->process(<span class="stringliteral">"#name_buf_value[lid] = #name_acc_value;"</span>) << std::endl;
<a name="l00278"></a>00278 stream << exprs[k]->process(<span class="stringliteral">"#name_buf[lid] = #name_acc;"</span>) << std::endl;
<a name="l00279"></a>00279 }
<a name="l00280"></a>00280
<a name="l00281"></a>00281
<a name="l00282"></a>00282 <span class="comment">//Reduce and write final result</span>
<a name="l00283"></a>00283 reduce_1d_local_memory(stream, <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>, exprs, <span class="stringliteral">"#name_buf"</span>, <span class="stringliteral">"#name_buf_value"</span>);
<a name="l00284"></a>00284
<a name="l00285"></a>00285 stream << <span class="stringliteral">"if (lid==0)"</span> << std::endl;
<a name="l00286"></a>00286 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00287"></a>00287 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00288"></a>00288 std::map<std::string, std::string> accessors;
<a name="l00289"></a>00289 accessors[<span class="stringliteral">"scalar_reduction"</span>] = <span class="stringliteral">"#name_buf[0]"</span>;
<a name="l00290"></a>00290 accessors[<span class="stringliteral">"scalar"</span>] = <span class="stringliteral">"*#pointer"</span>;
<a name="l00291"></a>00291 accessors[<span class="stringliteral">"vector"</span>] = <span class="stringliteral">"#pointer[#start]"</span>;
<a name="l00292"></a>00292 <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="l00293"></a>00293 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00294"></a>00294 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00295"></a>00295
<a name="l00296"></a>00296 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00297"></a>00297 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00298"></a>00298
<a name="l00299"></a>00299 <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="l00300"></a>00300 }
<a name="l00301"></a>00301
<a name="l00302"></a>00302 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="l00303"></a>00303 <span class="keyword"> </span>{
<a name="l00304"></a>00304 std::vector<std::string> result;
<a name="l00305"></a>00305 result.push_back(generate_impl(kernel_prefix + <span class="stringliteral">"_strided"</span>, statements, mappings, 1));
<a name="l00306"></a>00306 result.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>));
<a name="l00307"></a>00307 <span class="keywordflow">return</span> result;
<a name="l00308"></a>00308 }
<a name="l00309"></a>00309 <span class="keyword">public</span>:
<a name="l00310"></a><a class="code" href="classviennacl_1_1device__specific_1_1reduction__template.html#a387c3298d7df041316e29d0f22c0bea0">00310</a> <a class="code" href="classviennacl_1_1device__specific_1_1reduction__template.html#a387c3298d7df041316e29d0f22c0bea0">reduction_template</a>(<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html">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>, <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_1reduction__template.html">reduction_template</a>, <a class="code" href="structviennacl_1_1device__specific_1_1reduction__parameters.html">reduction_parameters</a>>(parameters, binding_policy) { }
<a name="l00311"></a>00311
<a name="l00312"></a><a class="code" href="classviennacl_1_1device__specific_1_1reduction__template.html#a259dc65aea9ba0534c06eee6d340e71a">00312</a> <span class="keywordtype">void</span> <a class="code" href="classviennacl_1_1device__specific_1_1reduction__template.html#a259dc65aea9ba0534c06eee6d340e71a">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="l00313"></a>00313 {
<a name="l00314"></a>00314 std::vector<scheduler::statement_node const *> reductions;
<a name="l00315"></a>00315 cl_uint size = 0;
<a name="l00316"></a>00316 <span class="keywordflow">for</span> (statements_container::data_type::const_iterator it = statements.<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html#a186e8a37443fdbc8deca9fc3f5fc64d5">data</a>().begin(); it != statements.<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html#a186e8a37443fdbc8deca9fc3f5fc64d5">data</a>().end(); ++it)
<a name="l00317"></a>00317 {
<a name="l00318"></a>00318 std::vector<vcl_size_t> reductions_idx;
<a name="l00319"></a>00319 <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>(*it, it->root(), <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>, reductions_idx), <span class="keyword">false</span>);
<a name="l00320"></a>00320 size = <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>(it->array(), reductions_idx[0]), <span class="keyword">false</span>));
<a name="l00321"></a>00321 <span class="keywordflow">for</span> (std::vector<vcl_size_t>::iterator itt = reductions_idx.begin(); itt != reductions_idx.end(); ++itt)
<a name="l00322"></a>00322 reductions.push_back(&it->array()[*itt]);
<a name="l00323"></a>00323 }
<a name="l00324"></a>00324
<a name="l00325"></a>00325 <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="l00326"></a>00326 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> scalartype_size = <a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#ad6565495d155b260e5c9f036c92aa8ae">utils::size_of</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>()).lhs.numeric_type);
<a name="l00327"></a>00327
<a name="l00328"></a>00328 <a class="code" href="classviennacl_1_1ocl_1_1kernel.html" title="Represents an OpenCL kernel within ViennaCL.">viennacl::ocl::kernel</a> * kernels[2];
<a name="l00329"></a>00329 <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="l00330"></a>00330 {
<a name="l00331"></a>00331 kernels[0] = &programs[0].program().get_kernel(kernel_prefix+<span class="stringliteral">"_strided_0"</span>);
<a name="l00332"></a>00332 kernels[1] = &programs[0].program().get_kernel(kernel_prefix+<span class="stringliteral">"_strided_1"</span>);
<a name="l00333"></a>00333 }
<a name="l00334"></a>00334 <span class="keywordflow">else</span>
<a name="l00335"></a>00335 {
<a name="l00336"></a>00336 kernels[0] = &programs[1].program().get_kernel(kernel_prefix+<span class="stringliteral">"_0"</span>);
<a name="l00337"></a>00337 kernels[1] = &programs[1].program().get_kernel(kernel_prefix+<span class="stringliteral">"_1"</span>);
<a name="l00338"></a>00338 }
<a name="l00339"></a>00339
<a name="l00340"></a>00340 kernels[0]-><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="l00341"></a>00341 kernels[0]-><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="l00342"></a>00342
<a name="l00343"></a>00343 kernels[1]-><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="l00344"></a>00344 kernels[1]-><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 name="l00345"></a>00345
<a name="l00346"></a>00346 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < 2; k++)
<a name="l00347"></a>00347 {
<a name="l00348"></a>00348 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> n_arg = 0;
<a name="l00349"></a>00349 kernels[k]-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#a86bff03b705c7cd32b1374b98477133e" title="Sets a char argument at the provided position.">arg</a>(n_arg++, size);
<a name="l00350"></a>00350 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> i = 0;
<a name="l00351"></a>00351 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> j = 0;
<a name="l00352"></a>00352 <span class="keywordflow">for</span> (std::vector<scheduler::statement_node const *>::const_iterator it = reductions.begin(); it != reductions.end(); ++it)
<a name="l00353"></a>00353 {
<a name="l00354"></a>00354 <span class="keywordflow">if</span> (<a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#ae0bee04daa630b804e130cfeecbbeb70">utils::is_index_reduction</a>((*it)->op))
<a name="l00355"></a>00355 {
<a name="l00356"></a>00356 <span class="keywordflow">if</span> (tmpidx_.size() <= j)
<a name="l00357"></a>00357 tmpidx_.push_back(kernels[k]-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#af60c5f1588e699e9a93a6fe879a45e76">context</a>().<a class="code" href="classviennacl_1_1ocl_1_1context.html#ac70d01bc7f8bd1f36c10eae811672d79" title="Creates a memory buffer within the context.">create_memory</a>(CL_MEM_READ_WRITE, <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.num_groups*4));
<a name="l00358"></a>00358 kernels[k]-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#a86bff03b705c7cd32b1374b98477133e" title="Sets a char argument at the provided position.">arg</a>(n_arg++, tmpidx_[j]);
<a name="l00359"></a>00359 j++;
<a name="l00360"></a>00360 }
<a name="l00361"></a>00361
<a name="l00362"></a>00362 <span class="keywordflow">if</span> (tmp_.size() <= i)
<a name="l00363"></a>00363 tmp_.push_back(kernels[k]-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#af60c5f1588e699e9a93a6fe879a45e76">context</a>().<a class="code" href="classviennacl_1_1ocl_1_1context.html#ac70d01bc7f8bd1f36c10eae811672d79" title="Creates a memory buffer within the context.">create_memory</a>(CL_MEM_READ_WRITE, <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.num_groups*scalartype_size));
<a name="l00364"></a>00364 kernels[k]-><a class="code" href="classviennacl_1_1ocl_1_1kernel.html#a86bff03b705c7cd32b1374b98477133e" title="Sets a char argument at the provided position.">arg</a>(n_arg++, tmp_[i]);
<a name="l00365"></a>00365 i++;
<a name="l00366"></a>00366 }
<a name="l00367"></a>00367 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a313b8339e22e2a1fdd5c2c885619f1ef">set_arguments</a>(statements, *kernels[k], n_arg);
<a name="l00368"></a>00368 }
<a name="l00369"></a>00369
<a name="l00370"></a>00370 <span class="keywordflow">for</span> (<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> k = 0; k < 2; k++)
<a name="l00371"></a>00371 <a class="code" href="namespaceviennacl_1_1ocl.html#a5f2022f653ea1cf364d20e3ff84dcada" title="Enqueues a kernel in the provided queue.">viennacl::ocl::enqueue</a>(*kernels[k]);
<a name="l00372"></a>00372
<a name="l00373"></a>00373 }
<a name="l00374"></a>00374
<a name="l00375"></a>00375 <span class="keyword">private</span>:
<a name="l00376"></a>00376 std::vector< viennacl::ocl::handle<cl_mem> > tmp_;
<a name="l00377"></a>00377 std::vector< viennacl::ocl::handle<cl_mem> > tmpidx_;
<a name="l00378"></a>00378 };
<a name="l00379"></a>00379
<a name="l00380"></a>00380 }
<a name="l00381"></a>00381 }
<a name="l00382"></a>00382
<a name="l00383"></a>00383 <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="reduction__template_8hpp.html">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>
|
