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
|
<!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/matrix_axpy_template.hpp Source File</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
$(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
</head>
<body>
<div id="top"><!-- do not remove this div! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td style="padding-left: 0.5em;">
<div id="projectname">ViennaCL - The Vienna Computing Library
 <span id="projectnumber">1.7.0</span>
</div>
<div id="projectbrief">Free open-source GPU-accelerated linear algebra and solver library.</div>
</td>
<td> <div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- Generated by Doxygen 1.7.6.1 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
</div>
<div id="side-nav" class="ui-resizable side-nav-resizable">
<div id="nav-tree">
<div id="nav-tree-contents">
</div>
</div>
<div id="splitbar" style="-moz-user-select:none;"
class="ui-resizable-handle">
</div>
</div>
<script type="text/javascript">
initNavTree('matrix__axpy__template_8hpp.html','');
</script>
<div id="doc-content">
<div class="header">
<div class="headertitle">
<div class="title">viennacl/device_specific/templates/matrix_axpy_template.hpp</div> </div>
</div><!--header-->
<div class="contents">
<a href="matrix__axpy__template_8hpp.html">Go to the documentation of this file.</a><div class="fragment"><pre class="fragment"><a name="l00001"></a>00001 <span class="preprocessor">#ifndef VIENNACL_DEVICE_SPECIFIC_TEMPLATES_MATRIX_AXPY_HPP</span>
<a name="l00002"></a>00002 <span class="preprocessor"></span><span class="preprocessor">#define VIENNACL_DEVICE_SPECIFIC_TEMPLATES_MATRIX_AXPY_HPP</span>
<a name="l00003"></a>00003 <span class="preprocessor"></span>
<a name="l00004"></a>00004 <span class="comment">/* =========================================================================</span>
<a name="l00005"></a>00005 <span class="comment"> Copyright (c) 2010-2015, Institute for Microelectronics,</span>
<a name="l00006"></a>00006 <span class="comment"> Institute for Analysis and Scientific Computing,</span>
<a name="l00007"></a>00007 <span class="comment"> TU Wien.</span>
<a name="l00008"></a>00008 <span class="comment"> Portions of this software are copyright by UChicago Argonne, LLC.</span>
<a name="l00009"></a>00009 <span class="comment"></span>
<a name="l00010"></a>00010 <span class="comment"> -----------------</span>
<a name="l00011"></a>00011 <span class="comment"> ViennaCL - The Vienna Computing Library</span>
<a name="l00012"></a>00012 <span class="comment"> -----------------</span>
<a name="l00013"></a>00013 <span class="comment"></span>
<a name="l00014"></a>00014 <span class="comment"> Project Head: Karl Rupp rupp@iue.tuwien.ac.at</span>
<a name="l00015"></a>00015 <span class="comment"></span>
<a name="l00016"></a>00016 <span class="comment"> (A list of authors and contributors can be found in the manual)</span>
<a name="l00017"></a>00017 <span class="comment"></span>
<a name="l00018"></a>00018 <span class="comment"> License: MIT (X11), see file LICENSE in the base directory</span>
<a name="l00019"></a>00019 <span class="comment">============================================================================= */</span>
<a name="l00020"></a>00020
<a name="l00021"></a>00021
<a name="l00027"></a>00027 <span class="preprocessor">#include <vector></span>
<a name="l00028"></a>00028
<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
<a name="l00037"></a>00037 <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="l00038"></a>00038
<a name="l00039"></a>00039 <span class="keyword">namespace </span>viennacl
<a name="l00040"></a>00040 {
<a name="l00041"></a>00041 <span class="keyword">namespace </span>device_specific
<a name="l00042"></a>00042 {
<a name="l00043"></a>00043
<a name="l00044"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html">00044</a> <span class="keyword">class </span><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html">matrix_axpy_parameters_type</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="l00045"></a>00045 {
<a name="l00046"></a>00046 <span class="keyword">public</span>:
<a name="l00047"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#aa09905d23f6d94e73d3b06d27b7ce78d">00047</a> <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#aa09905d23f6d94e73d3b06d27b7ce78d">matrix_axpy_parameters_type</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> _local_size_0, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _local_size_1,
<a name="l00049"></a>00049 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _num_groups_0, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> _num_groups_1,
<a name="l00050"></a>00050 <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, _local_size_0, _local_size_1, 1), <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#aa96d9f2ac2ea1e330e6c5aa8ddce578e">num_groups_0</a>(_num_groups_0), <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#afa3c5d69fb3546ed8ee633b6ed1a442b">num_groups_1</a>(_num_groups_1), <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#a2e16f2f724b74614f60617b54ca2acac">fetching_policy</a>(_fetching_policy){ }
<a name="l00051"></a>00051
<a name="l00052"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#aa96d9f2ac2ea1e330e6c5aa8ddce578e">00052</a> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#aa96d9f2ac2ea1e330e6c5aa8ddce578e">num_groups_0</a>;
<a name="l00053"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#afa3c5d69fb3546ed8ee633b6ed1a442b">00053</a> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#afa3c5d69fb3546ed8ee633b6ed1a442b">num_groups_1</a>;
<a name="l00054"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#a2e16f2f724b74614f60617b54ca2acac">00054</a> <a class="code" href="namespaceviennacl_1_1device__specific.html#acc80228a2075912f1505b042b6a9ff7c">fetching_policy_type</a> <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html#a2e16f2f724b74614f60617b54ca2acac">fetching_policy</a>;
<a name="l00055"></a>00055 };
<a name="l00056"></a>00056
<a name="l00057"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html">00057</a> <span class="keyword">class </span><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html">matrix_axpy_template</a> : <span class="keyword">public</span> <a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html">template_base_impl</a><matrix_axpy_template, matrix_axpy_parameters_type>
<a name="l00058"></a>00058 {
<a name="l00059"></a>00059 <span class="keyword">private</span>:
<a name="l00060"></a>00060 <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="l00061"></a>00061 <span class="keyword"> </span>{
<a name="l00062"></a>00062 <span class="keywordflow">if</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#af859aa51f8a2c20e654f63fdb34bb5b6">simd_width</a>>1)
<a name="l00063"></a>00063 <span class="keywordflow">return</span> TEMPLATE_INVALID_SIMD_WIDTH;
<a name="l00064"></a>00064 <span class="keywordflow">return</span> TEMPLATE_VALID;
<a name="l00065"></a>00065 }
<a name="l00066"></a>00066
<a name="l00067"></a>00067 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="l00068"></a>00068 <span class="keyword"> </span>{
<a name="l00069"></a>00069 std::string process_str;
<a name="l00070"></a>00070 <a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html">utils::kernel_generation_stream</a> stream;
<a name="l00071"></a>00071
<a name="l00072"></a>00072 std::string init0, upper_bound0, inc0, init1, upper_bound1, inc1;
<a name="l00073"></a>00073
<a name="l00074"></a>00074 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="l00075"></a>00075 <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="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
<a name="l00079"></a>00079 <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="l00080"></a>00080 <a class="code" href="namespaceviennacl_1_1device__specific_1_1tree__parsing.html#a7935d6fe2cb68ce032c9a177ed7f1d41">tree_parsing::process</a>(stream, <a class="code" href="namespaceviennacl_1_1device__specific.html#ae3c2256f348f5311933d23d2f699fa12a6f3d3028017d020e2fc0e4157449caff">PARENT_NODE_TYPE</a>, <span class="stringliteral">"matrix"</span>, <span class="stringliteral">"#pointer += $OFFSET{#start1, #start2};"</span>, statements, mappings);
<a name="l00081"></a>00081 <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="l00082"></a>00082
<a name="l00083"></a>00083 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#aef49b6bafebc3b3bdbe8814098912ea4">fetching_loop_info</a>(<a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.fetching_policy, <span class="stringliteral">"M"</span>, stream, init0, upper_bound0, inc0, <span class="stringliteral">"get_global_id(0)"</span>, <span class="stringliteral">"get_global_size(0)"</span>);
<a name="l00084"></a>00084 stream << <span class="stringliteral">"for(unsigned int i = "</span> << init0 << <span class="stringliteral">"; i < "</span> << upper_bound0 << <span class="stringliteral">"; i += "</span> << inc0 << <span class="stringliteral">")"</span> << std::endl;
<a name="l00085"></a>00085 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00086"></a>00086 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00087"></a>00087 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#aef49b6bafebc3b3bdbe8814098912ea4">fetching_loop_info</a>(<a class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.fetching_policy, <span class="stringliteral">"N"</span>, stream, init1, upper_bound1, inc1, <span class="stringliteral">"get_global_id(1)"</span>, <span class="stringliteral">"get_global_size(1)"</span>);
<a name="l00088"></a>00088 stream << <span class="stringliteral">"for(unsigned int j = "</span> << init1 << <span class="stringliteral">"; j < "</span> << upper_bound1 << <span class="stringliteral">"; j += "</span> << inc1 << <span class="stringliteral">")"</span> << std::endl;
<a name="l00089"></a>00089 stream << <span class="stringliteral">"{"</span> << std::endl;
<a name="l00090"></a>00090 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#aa2a76505b4c45bda548d59feca925d1a">inc_tab</a>();
<a name="l00091"></a>00091
<a name="l00092"></a>00092 process_str = <a class="code" href="namespaceviennacl_1_1device__specific_1_1utils.html#a18ef61de8cf8ec837e0107cc26f46da9">utils::append_width</a>(<span class="stringliteral">"#scalartype"</span>,simd_width) + <span class="stringliteral">" #namereg = "</span> + <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a80aea1b0b07097be9834c54daf0ca765">vload</a>(simd_width, <span class="stringliteral">"$OFFSET{i*#stride1,j*#stride2}"</span>, <span class="stringliteral">"#pointer"</span>)+ <span class="stringliteral">";"</span>;
<a name="l00093"></a>00093 <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>, process_str, statements, mappings);
<a name="l00094"></a>00094 <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_diag"</span>, <span class="stringliteral">"#scalartype #namereg = ((i + ((#diag_offset<0)?#diag_offset:0))!=(j-((#diag_offset>0)?#diag_offset:0)))?0:#pointer[min(i*#stride, j*#stride)];"</span>, statements, mappings);
<a name="l00095"></a>00095
<a name="l00096"></a>00096
<a name="l00097"></a>00097 std::map<std::string, std::string> accessors;
<a name="l00098"></a>00098 accessors[<span class="stringliteral">"matrix"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00099"></a>00099 accessors[<span class="stringliteral">"vector_diag"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00100"></a>00100 accessors[<span class="stringliteral">"scalar"</span>] = <span class="stringliteral">"#namereg"</span>;
<a name="l00101"></a>00101 <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="l00102"></a>00102
<a name="l00103"></a>00103 process_str = <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#aab0c82a5d605a94281f39b69a834f1b5">vstore</a>(simd_width, <span class="stringliteral">"#namereg"</span>, <span class="stringliteral">"$OFFSET{i*#stride1,j*#stride2}"</span>, <span class="stringliteral">"#pointer"</span>)+<span class="stringliteral">";"</span>;
<a name="l00104"></a>00104 <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#ae3c2256f348f5311933d23d2f699fa12afbf0ae324c55a1c703b3cc3e425a0bd6">LHS_NODE_TYPE</a>, <span class="stringliteral">"matrix"</span>, process_str, statements, mappings);
<a name="l00105"></a>00105
<a name="l00106"></a>00106 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00107"></a>00107 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00108"></a>00108 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00109"></a>00109 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00110"></a>00110
<a name="l00111"></a>00111 stream.<a class="code" href="classviennacl_1_1device__specific_1_1utils_1_1kernel__generation__stream.html#a39fa6c2e302d8c2163ca2340d9ccf1b9">dec_tab</a>();
<a name="l00112"></a>00112 stream << <span class="stringliteral">"}"</span> << std::endl;
<a name="l00113"></a>00113
<a name="l00114"></a>00114 <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="l00115"></a>00115 }
<a name="l00116"></a>00116
<a name="l00117"></a>00117 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="l00118"></a>00118 <span class="keyword"> </span>{
<a name="l00119"></a>00119 std::vector<std::string> res;
<a name="l00120"></a>00120 res.push_back(generate_impl(kernel_prefix, statements, mappings, 1));
<a name="l00121"></a>00121 <span class="keywordflow">return</span> res;
<a name="l00122"></a>00122 }
<a name="l00123"></a>00123
<a name="l00124"></a>00124 <span class="keyword">public</span>:
<a name="l00125"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html#aea91e39d4f34f999adc215875681f2d1">00125</a> <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html#aea91e39d4f34f999adc215875681f2d1">matrix_axpy_template</a>(<a class="code" href="structviennacl_1_1device__specific_1_1template__base_1_1parameters__type.html">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_1matrix__axpy__template.html">matrix_axpy_template</a>, <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__parameters__type.html">matrix_axpy_parameters_type</a>>(parameters, binding_policy), up_to_internal_size_(false){ }
<a name="l00126"></a>00126
<a name="l00127"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html#a3b489a3a6960699bfb85ffae95273882">00127</a> <span class="keywordtype">void</span> <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html#a3b489a3a6960699bfb85ffae95273882">up_to_internal_size</a>(<span class="keywordtype">bool</span> v)
<a name="l00128"></a>00128 {
<a name="l00129"></a>00129 up_to_internal_size_ = v;
<a name="l00130"></a>00130 }
<a name="l00131"></a>00131
<a name="l00132"></a><a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html#a12d1fe1746cfbde6fff6163b0cf15c75">00132</a> <span class="keywordtype">void</span> <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html#a12d1fe1746cfbde6fff6163b0cf15c75">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="l00133"></a>00133 {
<a name="l00134"></a>00134 <a class="code" href="classviennacl_1_1ocl_1_1kernel.html" title="Represents an OpenCL kernel within ViennaCL.">viennacl::ocl::kernel</a> & kernel = programs[0].program().get_kernel(kernel_prefix);
<a name="l00135"></a>00135
<a name="l00136"></a>00136 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="l00137"></a>00137 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="l00138"></a>00138 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="l00139"></a>00139 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 class="code" href="classviennacl_1_1device__specific_1_1template__base__impl.html#ac1a2b570d477d82f9cbd3d88dcb6e8c2">p_</a>.num_groups_1);
<a name="l00140"></a>00140
<a name="l00141"></a>00141 <a class="code" href="structviennacl_1_1scheduler_1_1statement__node.html" title="Main datastructure for an node in the statement tree.">scheduler::statement_node</a> <span class="keyword">const</span> & root = statements.<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html#a186e8a37443fdbc8deca9fc3f5fc64d5">data</a>().front().array()[statements.<a class="code" href="classviennacl_1_1device__specific_1_1statements__container.html#a186e8a37443fdbc8deca9fc3f5fc64d5">data</a>().front().root()];
<a name="l00142"></a>00142 <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> current_arg = 0;
<a name="l00143"></a>00143 <span class="keywordflow">if</span> (up_to_internal_size_)
<a name="l00144"></a>00144 {
<a name="l00145"></a>00145 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(root.<a class="code" href="structviennacl_1_1scheduler_1_1statement__node.html#a39cd9ba06df3613e4cc122a4b9d288a3">lhs</a>, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1internal__size1__fun.html">utils::internal_size1_fun</a>())));
<a name="l00146"></a>00146 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(root.<a class="code" href="structviennacl_1_1scheduler_1_1statement__node.html#a39cd9ba06df3613e4cc122a4b9d288a3">lhs</a>, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1internal__size2__fun.html">utils::internal_size2_fun</a>())));
<a name="l00147"></a>00147 }
<a name="l00148"></a>00148 <span class="keywordflow">else</span>
<a name="l00149"></a>00149 {
<a name="l00150"></a>00150 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(root.<a class="code" href="structviennacl_1_1scheduler_1_1statement__node.html#a39cd9ba06df3613e4cc122a4b9d288a3">lhs</a>, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1size1__fun.html">utils::size1_fun</a>())));
<a name="l00151"></a>00151 kernel.<a class="code" href="classviennacl_1_1ocl_1_1kernel.html#a86bff03b705c7cd32b1374b98477133e" title="Sets a char argument at the provided position.">arg</a>(current_arg++, cl_uint(utils::call_on_matrix(root.<a class="code" href="structviennacl_1_1scheduler_1_1statement__node.html#a39cd9ba06df3613e4cc122a4b9d288a3">lhs</a>, <a class="code" href="structviennacl_1_1device__specific_1_1utils_1_1size2__fun.html">utils::size2_fun</a>())));
<a name="l00152"></a>00152 }
<a name="l00153"></a>00153
<a name="l00154"></a>00154 <a class="code" href="classviennacl_1_1device__specific_1_1template__base.html#a313b8339e22e2a1fdd5c2c885619f1ef">set_arguments</a>(statements, kernel, current_arg);
<a name="l00155"></a>00155
<a name="l00156"></a>00156 <a class="code" href="classviennacl_1_1device__specific_1_1matrix__axpy__template.html#a12d1fe1746cfbde6fff6163b0cf15c75">viennacl::ocl::enqueue</a>(kernel);
<a name="l00157"></a>00157 }
<a name="l00158"></a>00158
<a name="l00159"></a>00159
<a name="l00160"></a>00160 <span class="keyword">private</span>:
<a name="l00161"></a>00161 <span class="keywordtype">bool</span> up_to_internal_size_;
<a name="l00162"></a>00162 };
<a name="l00163"></a>00163
<a name="l00164"></a>00164 }
<a name="l00165"></a>00165 }
<a name="l00166"></a>00166
<a name="l00167"></a>00167 <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="matrix__axpy__template_8hpp.html">matrix_axpy_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>
|
