Line data Source code
1 : /* -*- Mode: C++; tab-width: 4; indent-tabs-mode: nil; c-basic-offset: 4 -*- */
2 : /*
3 : * This file is part of the LibreOffice project.
4 : *
5 : * This Source Code Form is subject to the terms of the Mozilla Public
6 : * License, v. 2.0. If a copy of the MPL was not distributed with this
7 : * file, You can obtain one at http://mozilla.org/MPL/2.0/.
8 : */
9 :
10 : #include "opbase.hxx"
11 :
12 : using namespace formula;
13 :
14 : namespace sc { namespace opencl {
15 :
16 0 : DynamicKernelArgument::DynamicKernelArgument(const std::string &s,
17 : FormulaTreeNodeRef ft):
18 0 : mSymName(s), mFormulaTree(ft) {}
19 :
20 : /// Generate use/references to the argument
21 0 : void DynamicKernelArgument::GenDeclRef(std::stringstream &ss) const
22 : {
23 0 : ss << mSymName;
24 0 : }
25 :
26 0 : FormulaToken* DynamicKernelArgument::GetFormulaToken(void) const
27 : {
28 0 : return mFormulaTree->GetFormulaToken();
29 : }
30 :
31 0 : VectorRef::VectorRef(const std::string &s, FormulaTreeNodeRef ft, int idx):
32 0 : DynamicKernelArgument(s, ft), mpClmem(NULL), mnIndex(idx)
33 : {
34 0 : if (mnIndex)
35 : {
36 0 : std::stringstream ss;
37 0 : ss << mSymName << "s" << mnIndex;
38 0 : mSymName = ss.str();
39 : }
40 0 : }
41 :
42 0 : VectorRef::~VectorRef()
43 : {
44 0 : if (mpClmem) {
45 0 : cl_int ret = clReleaseMemObject(mpClmem);
46 0 : if (ret != CL_SUCCESS)
47 0 : throw OpenCLError(ret, __FILE__, __LINE__);
48 : }
49 0 : }
50 :
51 : /// Generate declaration
52 0 : void VectorRef::GenDecl(std::stringstream &ss) const
53 : {
54 0 : ss << "__global double *"<<mSymName;
55 0 : }
56 :
57 : /// When declared as input to a sliding window function
58 0 : void VectorRef::GenSlidingWindowDecl(std::stringstream &ss) const
59 : {
60 0 : VectorRef::GenDecl(ss);
61 0 : }
62 :
63 : /// When referenced in a sliding window function
64 0 : std::string VectorRef::GenSlidingWindowDeclRef(bool) const
65 : {
66 0 : std::stringstream ss;
67 : formula::SingleVectorRefToken *pSVR =
68 0 : dynamic_cast<formula::SingleVectorRefToken*>(DynamicKernelArgument::GetFormulaToken());
69 0 : if (pSVR)
70 0 : ss << "(gid0 < " << pSVR->GetArrayLength() << "?";
71 0 : ss << mSymName << "[gid0]";
72 0 : if (pSVR)
73 0 : ss << ":NAN)";
74 0 : return ss.str();
75 : }
76 :
77 0 : size_t VectorRef::GetWindowSize(void) const
78 : {
79 0 : FormulaToken *pCur = mFormulaTree->GetFormulaToken();
80 : assert(pCur);
81 0 : if (const formula::DoubleVectorRefToken* pCurDVR =
82 0 : dynamic_cast<const formula::DoubleVectorRefToken *>(pCur))
83 : {
84 0 : return pCurDVR->GetRefRowSize();
85 : }
86 0 : else if (dynamic_cast<const formula::SingleVectorRefToken *>(pCur))
87 : {
88 : // Prepare intermediate results (on CPU for now)
89 0 : return 1;
90 : }
91 : else
92 : {
93 0 : throw Unhandled();
94 : }
95 : }
96 :
97 0 : void Normal::GenSlidingWindowFunction(
98 : std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments)
99 : {
100 0 : ArgVector argVector;
101 0 : ss << "\ndouble " << sSymName;
102 0 : ss << "_"<< BinFuncName() <<"(";
103 0 : for (unsigned i = 0; i < vSubArguments.size(); i++)
104 : {
105 0 : if (i)
106 0 : ss << ",";
107 0 : vSubArguments[i]->GenSlidingWindowDecl(ss);
108 0 : argVector.push_back(vSubArguments[i]->GenSlidingWindowDeclRef());
109 : }
110 0 : ss << ") {\n\t";
111 0 : ss << "double tmp = " << GetBottom() <<";\n\t";
112 0 : ss << "int gid0 = get_global_id(0);\n\t";
113 0 : ss << "tmp = ";
114 0 : ss << Gen(argVector);
115 0 : ss << ";\n\t";
116 0 : ss << "return tmp;\n";
117 0 : ss << "}";
118 0 : }
119 :
120 0 : void CheckVariables::GenTmpVariables(
121 : std::stringstream & ss, SubArguments & vSubArguments)
122 : {
123 0 : for(unsigned i=0;i<vSubArguments.size();i++)
124 : {
125 0 : ss << " double tmp";
126 0 : ss << i;
127 0 : ss <<";\n";
128 : }
129 0 : }
130 :
131 0 : void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss,
132 : SubArguments &vSubArguments, int argumentNum)
133 : {
134 0 : int i = argumentNum;
135 : #ifdef ISNAN
136 0 : if(vSubArguments[i]->GetFormulaToken()->GetType() ==
137 : formula::svSingleVectorRef)
138 : {
139 : const formula::SingleVectorRefToken*pTmpDVR1= static_cast<const
140 0 : formula::SingleVectorRefToken *>(vSubArguments[i]->GetFormulaToken());
141 0 : ss<< " if(singleIndex>=";
142 0 : ss<< pTmpDVR1->GetArrayLength();
143 0 : ss<<" ||";
144 : }
145 0 : if(vSubArguments[i]->GetFormulaToken()->GetType() ==
146 : formula::svDoubleVectorRef)
147 : {
148 : const formula::DoubleVectorRefToken*pTmpDVR2= static_cast<const
149 0 : formula::DoubleVectorRefToken *>(vSubArguments[i]->GetFormulaToken());
150 0 : ss<< " if(doubleIndex>=";
151 0 : ss<< pTmpDVR2->GetArrayLength();
152 0 : ss<<" ||";
153 : }
154 0 : if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble ||
155 0 : vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush)
156 : {
157 0 : ss<< " if(";
158 : }
159 0 : ss<< "isNan(";
160 0 : ss<< vSubArguments[i]->GenSlidingWindowDeclRef();
161 0 : ss<<"))\n";
162 0 : ss<< " tmp";
163 0 : ss<< i;
164 0 : ss <<"=0;\n else \n";
165 : #endif
166 0 : ss <<" tmp";
167 0 : ss <<i;
168 0 : ss << "=";
169 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef();
170 0 : ss<<";\n";
171 0 : }
172 :
173 0 : void CheckVariables::CheckSubArgumentIsNan2( std::stringstream & ss,
174 : SubArguments &vSubArguments, int argumentNum, std::string p)
175 : {
176 0 : int i = argumentNum;
177 0 : if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble)
178 : {
179 0 : ss <<" tmp";
180 0 : ss <<i;
181 0 : ss << "=";
182 0 : vSubArguments[i]->GenDeclRef(ss);
183 0 : ss<<";\n";
184 0 : return;
185 : }
186 :
187 : #ifdef ISNAN
188 0 : ss<< " tmp";
189 0 : ss<< i;
190 0 : ss<< "= fsum(";
191 0 : vSubArguments[i]->GenDeclRef(ss);
192 0 : if(vSubArguments[i]->GetFormulaToken()->GetType() ==
193 : formula::svDoubleVectorRef)
194 0 : ss<<"["<< p.c_str()<< "]";
195 0 : else if(vSubArguments[i]->GetFormulaToken()->GetType() ==
196 : formula::svSingleVectorRef)
197 0 : ss<<"[get_group_id(1)]";
198 0 : ss<<", 0);\n";
199 : #else
200 : ss <<" tmp";
201 : ss <<i;
202 : ss << "=";
203 : vSubArguments[i]->GenDeclRef(ss);
204 : if(vSubArguments[i]->GetFormulaToken()->GetType() ==
205 : formula::svDoubleVectorRef)
206 : ss<<"["<< p.c_str()<< "]";
207 : else if(vSubArguments[i]->GetFormulaToken()->GetType() ==
208 : formula::svSingleVectorRef)
209 : ss<<"[get_group_id(1)]";
210 :
211 : ss<<";\n";
212 : #endif
213 : }
214 :
215 0 : void CheckVariables::CheckAllSubArgumentIsNan(
216 : std::stringstream & ss, SubArguments & vSubArguments)
217 : {
218 0 : ss<<" int k = gid0;\n";
219 0 : for(unsigned i=0;i<vSubArguments.size();i++)
220 : {
221 0 : CheckSubArgumentIsNan(ss,vSubArguments,i);
222 : }
223 0 : }
224 :
225 0 : void CheckVariables::UnrollDoubleVector( std::stringstream & ss,
226 : std::stringstream & unrollstr,const formula::DoubleVectorRefToken* pCurDVR,
227 : int nCurWindowSize)
228 : {
229 0 : int unrollSize = 16;
230 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) {
231 0 : ss << " loop = ("<<nCurWindowSize<<" - gid0)/";
232 0 : ss << unrollSize<<";\n";
233 0 : } else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) {
234 0 : ss << " loop = ("<<nCurWindowSize<<" + gid0)/";
235 0 : ss << unrollSize<<";\n";
236 :
237 : } else {
238 0 : ss << " loop = "<<nCurWindowSize<<"/"<< unrollSize<<";\n";
239 : }
240 :
241 0 : ss << " for ( int j = 0;j< loop; j++)\n";
242 0 : ss << " {\n";
243 0 : ss << " int i = ";
244 0 : if (!pCurDVR->IsStartFixed()&& pCurDVR->IsEndFixed()) {
245 0 : ss << "gid0 + j * "<< unrollSize <<";\n";
246 : }else {
247 0 : ss << "j * "<< unrollSize <<";\n";
248 : }
249 :
250 0 : if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
251 : {
252 0 : ss << " int doubleIndex = i+gid0;\n";
253 : }else
254 : {
255 0 : ss << " int doubleIndex = i;\n";
256 : }
257 :
258 0 : for(int j =0;j < unrollSize;j++)
259 : {
260 0 : ss << unrollstr.str();
261 0 : ss << "i++;\n";
262 0 : ss << "doubleIndex++;\n";
263 : }
264 0 : ss << " }\n";
265 0 : ss << " for (int i = ";
266 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) {
267 0 : ss << "gid0 + loop *"<<unrollSize<<"; i < ";
268 0 : ss << nCurWindowSize <<"; i++)\n";
269 0 : } else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) {
270 0 : ss << "0 + loop *"<<unrollSize<<"; i < gid0+";
271 0 : ss << nCurWindowSize <<"; i++)\n";
272 : } else {
273 0 : ss << "0 + loop *"<<unrollSize<<"; i < ";
274 0 : ss << nCurWindowSize <<"; i++)\n";
275 : }
276 0 : ss << " {\n";
277 0 : if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
278 : {
279 0 : ss << " int doubleIndex = i+gid0;\n";
280 : }else
281 : {
282 0 : ss << " int doubleIndex = i;\n";
283 : }
284 0 : ss << unrollstr.str();
285 0 : ss << " }\n";
286 0 : }
287 :
288 : }}
289 :
290 : /* vim:set shiftwidth=4 softtabstop=4 expandtab: */
|
