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 <opencl/openclwrapper.hxx>
11 :
12 : #include "opbase.hxx"
13 :
14 : using namespace formula;
15 :
16 : namespace sc { namespace opencl {
17 :
18 0 : UnhandledToken::UnhandledToken(
19 : formula::FormulaToken* t, const char* m, const std::string& fn, int ln ) :
20 0 : mToken(t), mMessage(m), mFile(fn), mLineNumber(ln) {}
21 :
22 0 : OpenCLError::OpenCLError( const std::string& function, cl_int error, const std::string& file, int line ) :
23 0 : mFunction(function), mError(error), mFile(file), mLineNumber(line)
24 : {
25 : // Not sure if this SAL_INFO() is useful; the place in
26 : // CLInterpreterContext::launchKernel() where OpenCLError is
27 : // caught already uses SAL_WARN() to display it.
28 :
29 : // SAL_INFO("sc.opencl", "OpenCL error: " << ::opencl::errorString(mError));
30 0 : }
31 :
32 0 : Unhandled::Unhandled( const std::string& fn, int ln ) :
33 0 : mFile(fn), mLineNumber(ln) {}
34 :
35 0 : DynamicKernelArgument::DynamicKernelArgument( const ScCalcConfig& config, const std::string& s,
36 : FormulaTreeNodeRef ft ) :
37 0 : mCalcConfig(config), mSymName(s), mFormulaTree(ft) { }
38 :
39 0 : std::string DynamicKernelArgument::GenDoubleSlidingWindowDeclRef( bool ) const
40 : {
41 0 : return std::string("");
42 : }
43 :
44 : /// When Mix, it will be called
45 0 : std::string DynamicKernelArgument::GenStringSlidingWindowDeclRef( bool ) const
46 : {
47 0 : return std::string("");
48 : }
49 :
50 0 : bool DynamicKernelArgument::IsMixedArgument() const
51 : {
52 0 : return false;
53 : }
54 :
55 : /// Generate use/references to the argument
56 0 : void DynamicKernelArgument::GenDeclRef( std::stringstream& ss ) const
57 : {
58 0 : ss << mSymName;
59 0 : }
60 :
61 0 : void DynamicKernelArgument::GenNumDeclRef( std::stringstream& ss ) const
62 : {
63 0 : ss << ",";
64 0 : }
65 :
66 0 : void DynamicKernelArgument::GenStringDeclRef( std::stringstream& ss ) const
67 : {
68 0 : ss << ",";
69 0 : }
70 :
71 0 : void DynamicKernelArgument::GenSlidingWindowFunction( std::stringstream& ) {}
72 :
73 0 : FormulaToken* DynamicKernelArgument::GetFormulaToken() const
74 : {
75 0 : return mFormulaTree->GetFormulaToken();
76 : }
77 :
78 0 : std::string DynamicKernelArgument::DumpOpName() const
79 : {
80 0 : return std::string("");
81 : }
82 :
83 0 : void DynamicKernelArgument::DumpInlineFun( std::set<std::string>&, std::set<std::string>& ) const {}
84 :
85 0 : const std::string& DynamicKernelArgument::GetName() const
86 : {
87 0 : return mSymName;
88 : }
89 :
90 0 : bool DynamicKernelArgument::NeedParallelReduction() const
91 : {
92 0 : return false;
93 : }
94 :
95 0 : VectorRef::VectorRef( const ScCalcConfig& config, const std::string& s, FormulaTreeNodeRef ft, int idx ) :
96 0 : DynamicKernelArgument(config, s, ft), mpClmem(NULL), mnIndex(idx)
97 : {
98 0 : if (mnIndex)
99 : {
100 0 : std::stringstream ss;
101 0 : ss << mSymName << "s" << mnIndex;
102 0 : mSymName = ss.str();
103 : }
104 0 : }
105 :
106 0 : VectorRef::~VectorRef()
107 : {
108 0 : if (mpClmem)
109 : {
110 : cl_int err;
111 0 : err = clReleaseMemObject(mpClmem);
112 : SAL_WARN_IF(err != CL_SUCCESS, "sc.opencl", "clReleaseMemObject failed: " << ::opencl::errorString(err));
113 : }
114 0 : }
115 :
116 : /// Generate declaration
117 0 : void VectorRef::GenDecl( std::stringstream& ss ) const
118 : {
119 0 : ss << "__global double *" << mSymName;
120 0 : }
121 :
122 : /// When declared as input to a sliding window function
123 0 : void VectorRef::GenSlidingWindowDecl( std::stringstream& ss ) const
124 : {
125 0 : VectorRef::GenDecl(ss);
126 0 : }
127 :
128 : /// When referenced in a sliding window function
129 0 : std::string VectorRef::GenSlidingWindowDeclRef( bool nested ) const
130 : {
131 0 : std::stringstream ss;
132 : formula::SingleVectorRefToken* pSVR =
133 0 : dynamic_cast<formula::SingleVectorRefToken*>(DynamicKernelArgument::GetFormulaToken());
134 0 : if (pSVR && !nested)
135 0 : ss << "(gid0 < " << pSVR->GetArrayLength() << "?";
136 0 : ss << mSymName << "[gid0]";
137 0 : if (pSVR && !nested)
138 0 : ss << ":NAN)";
139 0 : return ss.str();
140 : }
141 :
142 0 : void VectorRef::GenSlidingWindowFunction( std::stringstream& ) {}
143 :
144 0 : size_t VectorRef::GetWindowSize() const
145 : {
146 0 : FormulaToken* pCur = mFormulaTree->GetFormulaToken();
147 : assert(pCur);
148 0 : if (const formula::DoubleVectorRefToken* pCurDVR =
149 0 : dynamic_cast<const formula::DoubleVectorRefToken*>(pCur))
150 : {
151 0 : return pCurDVR->GetRefRowSize();
152 : }
153 0 : else if (dynamic_cast<const formula::SingleVectorRefToken*>(pCur))
154 : {
155 : // Prepare intermediate results (on CPU for now)
156 0 : return 1;
157 : }
158 : else
159 : {
160 0 : throw Unhandled();
161 : }
162 : }
163 :
164 0 : std::string VectorRef::DumpOpName() const
165 : {
166 0 : return std::string("");
167 : }
168 :
169 0 : void VectorRef::DumpInlineFun( std::set<std::string>&, std::set<std::string>& ) const {}
170 :
171 0 : const std::string& VectorRef::GetName() const
172 : {
173 0 : return mSymName;
174 : }
175 :
176 0 : cl_mem VectorRef::GetCLBuffer() const
177 : {
178 0 : return mpClmem;
179 : }
180 :
181 0 : bool VectorRef::NeedParallelReduction() const
182 : {
183 0 : return false;
184 : }
185 :
186 0 : void Normal::GenSlidingWindowFunction(
187 : std::stringstream& ss, const std::string& sSymName, SubArguments& vSubArguments )
188 : {
189 0 : ArgVector argVector;
190 0 : ss << "\ndouble " << sSymName;
191 0 : ss << "_" << BinFuncName() << "(";
192 0 : for (size_t i = 0; i < vSubArguments.size(); i++)
193 : {
194 0 : if (i)
195 0 : ss << ",";
196 0 : vSubArguments[i]->GenSlidingWindowDecl(ss);
197 0 : argVector.push_back(vSubArguments[i]->GenSlidingWindowDeclRef());
198 : }
199 0 : ss << ") {\n\t";
200 0 : ss << "double tmp = " << GetBottom() << ";\n\t";
201 0 : ss << "int gid0 = get_global_id(0);\n\t";
202 0 : ss << "tmp = ";
203 0 : ss << Gen(argVector);
204 0 : ss << ";\n\t";
205 0 : ss << "return tmp;\n";
206 0 : ss << "}";
207 0 : }
208 :
209 0 : void CheckVariables::GenTmpVariables(
210 : std::stringstream& ss, SubArguments& vSubArguments )
211 : {
212 0 : for (size_t i = 0; i < vSubArguments.size(); i++)
213 : {
214 0 : ss << " double tmp";
215 0 : ss << i;
216 0 : ss << ";\n";
217 : }
218 0 : }
219 :
220 0 : void CheckVariables::CheckSubArgumentIsNan( std::stringstream& ss,
221 : SubArguments& vSubArguments, int argumentNum )
222 : {
223 0 : int i = argumentNum;
224 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() ==
225 : formula::svSingleVectorRef)
226 : {
227 : const formula::SingleVectorRefToken* pTmpDVR1 =
228 0 : static_cast<const formula::SingleVectorRefToken*>(vSubArguments[i]->GetFormulaToken());
229 0 : ss << " if(singleIndex>=";
230 0 : ss << pTmpDVR1->GetArrayLength();
231 0 : ss << " ||";
232 0 : ss << "isNan(";
233 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(true);
234 0 : ss << "))\n";
235 0 : ss << " tmp";
236 0 : ss << i;
237 0 : ss << "=0;\n else \n";
238 0 : ss << " tmp";
239 0 : ss << i;
240 0 : ss << "=";
241 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(true);
242 0 : ss << ";\n";
243 : }
244 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() ==
245 : formula::svDoubleVectorRef)
246 : {
247 : const formula::DoubleVectorRefToken* pTmpDVR2 =
248 0 : static_cast<const formula::DoubleVectorRefToken*>(vSubArguments[i]->GetFormulaToken());
249 0 : ss << " if(doubleIndex>=";
250 0 : ss << pTmpDVR2->GetArrayLength();
251 0 : ss << " ||";
252 0 : ss << "isNan(";
253 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(false);
254 0 : ss << "))\n";
255 0 : ss << " tmp";
256 0 : ss << i;
257 0 : ss << "=0;\n else \n";
258 0 : ss << " tmp";
259 0 : ss << i;
260 0 : ss << "=";
261 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(false);
262 0 : ss << ";\n";
263 : }
264 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble ||
265 0 : vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush)
266 : {
267 0 : ss << " if(";
268 0 : ss << "isNan(";
269 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef();
270 0 : ss << "))\n";
271 0 : ss << " tmp";
272 0 : ss << i;
273 0 : ss << "=0;\n else \n";
274 0 : ss << " tmp";
275 0 : ss << i;
276 0 : ss << "=";
277 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef();
278 0 : ss << ";\n";
279 :
280 : }
281 :
282 0 : }
283 :
284 0 : void CheckVariables::CheckSubArgumentIsNan2( std::stringstream& ss,
285 : SubArguments& vSubArguments, int argumentNum, std::string p )
286 : {
287 0 : int i = argumentNum;
288 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble)
289 : {
290 0 : ss << " tmp";
291 0 : ss << i;
292 0 : ss << "=";
293 0 : vSubArguments[i]->GenDeclRef(ss);
294 0 : ss << ";\n";
295 0 : return;
296 : }
297 :
298 0 : ss << " tmp";
299 0 : ss << i;
300 0 : ss << "= fsum(";
301 0 : vSubArguments[i]->GenDeclRef(ss);
302 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() ==
303 : formula::svDoubleVectorRef)
304 0 : ss << "[" << p.c_str() << "]";
305 0 : else if (vSubArguments[i]->GetFormulaToken()->GetType() ==
306 : formula::svSingleVectorRef)
307 0 : ss << "[get_group_id(1)]";
308 0 : ss << ", 0);\n";
309 : }
310 :
311 0 : void CheckVariables::CheckAllSubArgumentIsNan(
312 : std::stringstream& ss, SubArguments& vSubArguments )
313 : {
314 0 : ss << " int k = gid0;\n";
315 0 : for (size_t i = 0; i < vSubArguments.size(); i++)
316 : {
317 0 : CheckSubArgumentIsNan(ss, vSubArguments, i);
318 : }
319 0 : }
320 :
321 0 : void CheckVariables::UnrollDoubleVector( std::stringstream& ss,
322 : std::stringstream& unrollstr, const formula::DoubleVectorRefToken* pCurDVR,
323 : int nCurWindowSize )
324 : {
325 0 : int unrollSize = 16;
326 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed())
327 : {
328 0 : ss << " loop = (" << nCurWindowSize << " - gid0)/";
329 0 : ss << unrollSize << ";\n";
330 : }
331 0 : else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
332 : {
333 0 : ss << " loop = (" << nCurWindowSize << " + gid0)/";
334 0 : ss << unrollSize << ";\n";
335 :
336 : }
337 : else
338 : {
339 0 : ss << " loop = " << nCurWindowSize << "/" << unrollSize << ";\n";
340 : }
341 :
342 0 : ss << " for ( int j = 0;j< loop; j++)\n";
343 0 : ss << " {\n";
344 0 : ss << " int i = ";
345 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed())
346 : {
347 0 : ss << "gid0 + j * " << unrollSize << ";\n";
348 : }
349 : else
350 : {
351 0 : ss << "j * " << unrollSize << ";\n";
352 : }
353 :
354 0 : if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
355 : {
356 0 : ss << " int doubleIndex = i+gid0;\n";
357 : }
358 : else
359 : {
360 0 : ss << " int doubleIndex = i;\n";
361 : }
362 :
363 0 : for (int j = 0; j < unrollSize; j++)
364 : {
365 0 : ss << unrollstr.str();
366 0 : ss << "i++;\n";
367 0 : ss << "doubleIndex++;\n";
368 : }
369 0 : ss << " }\n";
370 0 : ss << " for (int i = ";
371 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed())
372 : {
373 0 : ss << "gid0 + loop *" << unrollSize << "; i < ";
374 0 : ss << nCurWindowSize << "; i++)\n";
375 : }
376 0 : else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
377 : {
378 0 : ss << "0 + loop *" << unrollSize << "; i < gid0+";
379 0 : ss << nCurWindowSize << "; i++)\n";
380 : }
381 : else
382 : {
383 0 : ss << "0 + loop *" << unrollSize << "; i < ";
384 0 : ss << nCurWindowSize << "; i++)\n";
385 : }
386 0 : ss << " {\n";
387 0 : if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
388 : {
389 0 : ss << " int doubleIndex = i+gid0;\n";
390 : }
391 : else
392 : {
393 0 : ss << " int doubleIndex = i;\n";
394 : }
395 0 : ss << unrollstr.str();
396 0 : ss << " }\n";
397 0 : }
398 :
399 : }}
400 :
401 : /* vim:set shiftwidth=4 softtabstop=4 expandtab: */
|
