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 : UnhandledToken::UnhandledToken(
17 : formula::FormulaToken* t, const char* m, const std::string& fn, int ln ) :
18 0 : mToken(t), mMessage(m), mFile(fn), mLineNumber(ln) {}
19 :
20 0 : OpenCLError::OpenCLError( cl_int err, const std::string& fn, int ln ) :
21 0 : mError(err), mFile(fn), mLineNumber(ln)
22 : {
23 : SAL_INFO("sc.opencl", "OpenCLError:" << mError << ": " << strerror(mError));
24 0 : }
25 :
26 0 : const char* OpenCLError::strerror( cl_int i ) const
27 : {
28 : #define CASE(val) case val: return #val
29 0 : switch (i)
30 : {
31 0 : CASE(CL_SUCCESS);
32 0 : CASE(CL_DEVICE_NOT_FOUND);
33 0 : CASE(CL_DEVICE_NOT_AVAILABLE);
34 0 : CASE(CL_COMPILER_NOT_AVAILABLE);
35 0 : CASE(CL_MEM_OBJECT_ALLOCATION_FAILURE);
36 0 : CASE(CL_OUT_OF_RESOURCES);
37 0 : CASE(CL_OUT_OF_HOST_MEMORY);
38 0 : CASE(CL_PROFILING_INFO_NOT_AVAILABLE);
39 0 : CASE(CL_MEM_COPY_OVERLAP);
40 0 : CASE(CL_IMAGE_FORMAT_MISMATCH);
41 0 : CASE(CL_IMAGE_FORMAT_NOT_SUPPORTED);
42 0 : CASE(CL_BUILD_PROGRAM_FAILURE);
43 0 : CASE(CL_MAP_FAILURE);
44 0 : CASE(CL_INVALID_VALUE);
45 0 : CASE(CL_INVALID_DEVICE_TYPE);
46 0 : CASE(CL_INVALID_PLATFORM);
47 0 : CASE(CL_INVALID_DEVICE);
48 0 : CASE(CL_INVALID_CONTEXT);
49 0 : CASE(CL_INVALID_QUEUE_PROPERTIES);
50 0 : CASE(CL_INVALID_COMMAND_QUEUE);
51 0 : CASE(CL_INVALID_HOST_PTR);
52 0 : CASE(CL_INVALID_MEM_OBJECT);
53 0 : CASE(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
54 0 : CASE(CL_INVALID_IMAGE_SIZE);
55 0 : CASE(CL_INVALID_SAMPLER);
56 0 : CASE(CL_INVALID_BINARY);
57 0 : CASE(CL_INVALID_BUILD_OPTIONS);
58 0 : CASE(CL_INVALID_PROGRAM);
59 0 : CASE(CL_INVALID_PROGRAM_EXECUTABLE);
60 0 : CASE(CL_INVALID_KERNEL_NAME);
61 0 : CASE(CL_INVALID_KERNEL_DEFINITION);
62 0 : CASE(CL_INVALID_KERNEL);
63 0 : CASE(CL_INVALID_ARG_INDEX);
64 0 : CASE(CL_INVALID_ARG_VALUE);
65 0 : CASE(CL_INVALID_ARG_SIZE);
66 0 : CASE(CL_INVALID_KERNEL_ARGS);
67 0 : CASE(CL_INVALID_WORK_DIMENSION);
68 0 : CASE(CL_INVALID_WORK_GROUP_SIZE);
69 0 : CASE(CL_INVALID_WORK_ITEM_SIZE);
70 0 : CASE(CL_INVALID_GLOBAL_OFFSET);
71 0 : CASE(CL_INVALID_EVENT_WAIT_LIST);
72 0 : CASE(CL_INVALID_EVENT);
73 0 : CASE(CL_INVALID_OPERATION);
74 0 : CASE(CL_INVALID_GL_OBJECT);
75 0 : CASE(CL_INVALID_BUFFER_SIZE);
76 0 : CASE(CL_INVALID_MIP_LEVEL);
77 0 : CASE(CL_INVALID_GLOBAL_WORK_SIZE);
78 : default:
79 0 : return "Unknown OpenCL error code";
80 : }
81 : #undef CASE
82 : }
83 :
84 0 : Unhandled::Unhandled( const std::string& fn, int ln ) :
85 0 : mFile(fn), mLineNumber(ln) {}
86 :
87 0 : DynamicKernelArgument::DynamicKernelArgument( const std::string& s,
88 : FormulaTreeNodeRef ft ) :
89 0 : mSymName(s), mFormulaTree(ft) { }
90 :
91 0 : std::string DynamicKernelArgument::GenDoubleSlidingWindowDeclRef( bool ) const
92 : {
93 0 : return std::string("");
94 : }
95 :
96 : /// When Mix, it will be called
97 0 : std::string DynamicKernelArgument::GenStringSlidingWindowDeclRef( bool ) const
98 : {
99 0 : return std::string("");
100 : }
101 :
102 0 : bool DynamicKernelArgument::IsMixedArgument() const
103 : {
104 0 : return false;
105 : }
106 :
107 : /// Generate use/references to the argument
108 0 : void DynamicKernelArgument::GenDeclRef( std::stringstream& ss ) const
109 : {
110 0 : ss << mSymName;
111 0 : }
112 :
113 0 : void DynamicKernelArgument::GenNumDeclRef( std::stringstream& ss ) const
114 : {
115 0 : ss << ",";
116 0 : }
117 :
118 0 : void DynamicKernelArgument::GenStringDeclRef( std::stringstream& ss ) const
119 : {
120 0 : ss << ",";
121 0 : }
122 :
123 0 : void DynamicKernelArgument::GenSlidingWindowFunction( std::stringstream& ) {}
124 :
125 0 : FormulaToken* DynamicKernelArgument::GetFormulaToken() const
126 : {
127 0 : return mFormulaTree->GetFormulaToken();
128 : }
129 :
130 0 : std::string DynamicKernelArgument::DumpOpName() const
131 : {
132 0 : return std::string("");
133 : }
134 :
135 0 : void DynamicKernelArgument::DumpInlineFun( std::set<std::string>&, std::set<std::string>& ) const {}
136 :
137 0 : const std::string& DynamicKernelArgument::GetName() const
138 : {
139 0 : return mSymName;
140 : }
141 :
142 0 : bool DynamicKernelArgument::NeedParallelReduction() const
143 : {
144 0 : return false;
145 : }
146 :
147 0 : VectorRef::VectorRef( const std::string& s, FormulaTreeNodeRef ft, int idx ) :
148 0 : DynamicKernelArgument(s, ft), mpClmem(NULL), mnIndex(idx)
149 : {
150 0 : if (mnIndex)
151 : {
152 0 : std::stringstream ss;
153 0 : ss << mSymName << "s" << mnIndex;
154 0 : mSymName = ss.str();
155 : }
156 0 : }
157 :
158 0 : VectorRef::~VectorRef()
159 : {
160 0 : if (mpClmem)
161 : {
162 0 : clReleaseMemObject(mpClmem);
163 : }
164 0 : }
165 :
166 : /// Generate declaration
167 0 : void VectorRef::GenDecl( std::stringstream& ss ) const
168 : {
169 0 : ss << "__global double *" << mSymName;
170 0 : }
171 :
172 : /// When declared as input to a sliding window function
173 0 : void VectorRef::GenSlidingWindowDecl( std::stringstream& ss ) const
174 : {
175 0 : VectorRef::GenDecl(ss);
176 0 : }
177 :
178 : /// When referenced in a sliding window function
179 0 : std::string VectorRef::GenSlidingWindowDeclRef( bool nested ) const
180 : {
181 0 : std::stringstream ss;
182 : formula::SingleVectorRefToken* pSVR =
183 0 : dynamic_cast<formula::SingleVectorRefToken*>(DynamicKernelArgument::GetFormulaToken());
184 0 : if (pSVR && !nested)
185 0 : ss << "(gid0 < " << pSVR->GetArrayLength() << "?";
186 0 : ss << mSymName << "[gid0]";
187 0 : if (pSVR && !nested)
188 0 : ss << ":NAN)";
189 0 : return ss.str();
190 : }
191 :
192 0 : void VectorRef::GenSlidingWindowFunction( std::stringstream& ) {}
193 :
194 0 : size_t VectorRef::GetWindowSize() const
195 : {
196 0 : FormulaToken* pCur = mFormulaTree->GetFormulaToken();
197 : assert(pCur);
198 0 : if (const formula::DoubleVectorRefToken* pCurDVR =
199 0 : dynamic_cast<const formula::DoubleVectorRefToken*>(pCur))
200 : {
201 0 : return pCurDVR->GetRefRowSize();
202 : }
203 0 : else if (dynamic_cast<const formula::SingleVectorRefToken*>(pCur))
204 : {
205 : // Prepare intermediate results (on CPU for now)
206 0 : return 1;
207 : }
208 : else
209 : {
210 0 : throw Unhandled();
211 : }
212 : }
213 :
214 0 : std::string VectorRef::DumpOpName() const
215 : {
216 0 : return std::string("");
217 : }
218 :
219 0 : void VectorRef::DumpInlineFun( std::set<std::string>&, std::set<std::string>& ) const {}
220 :
221 0 : const std::string& VectorRef::GetName() const
222 : {
223 0 : return mSymName;
224 : }
225 :
226 0 : cl_mem VectorRef::GetCLBuffer() const
227 : {
228 0 : return mpClmem;
229 : }
230 :
231 0 : bool VectorRef::NeedParallelReduction() const
232 : {
233 0 : return false;
234 : }
235 :
236 0 : void Normal::GenSlidingWindowFunction(
237 : std::stringstream& ss, const std::string& sSymName, SubArguments& vSubArguments )
238 : {
239 0 : ArgVector argVector;
240 0 : ss << "\ndouble " << sSymName;
241 0 : ss << "_" << BinFuncName() << "(";
242 0 : for (unsigned i = 0; i < vSubArguments.size(); i++)
243 : {
244 0 : if (i)
245 0 : ss << ",";
246 0 : vSubArguments[i]->GenSlidingWindowDecl(ss);
247 0 : argVector.push_back(vSubArguments[i]->GenSlidingWindowDeclRef());
248 : }
249 0 : ss << ") {\n\t";
250 0 : ss << "double tmp = " << GetBottom() << ";\n\t";
251 0 : ss << "int gid0 = get_global_id(0);\n\t";
252 0 : ss << "tmp = ";
253 0 : ss << Gen(argVector);
254 0 : ss << ";\n\t";
255 0 : ss << "return tmp;\n";
256 0 : ss << "}";
257 0 : }
258 :
259 0 : void CheckVariables::GenTmpVariables(
260 : std::stringstream& ss, SubArguments& vSubArguments )
261 : {
262 0 : for (unsigned i = 0; i < vSubArguments.size(); i++)
263 : {
264 0 : ss << " double tmp";
265 0 : ss << i;
266 0 : ss << ";\n";
267 : }
268 0 : }
269 :
270 0 : void CheckVariables::CheckSubArgumentIsNan( std::stringstream& ss,
271 : SubArguments& vSubArguments, int argumentNum )
272 : {
273 0 : int i = argumentNum;
274 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() ==
275 : formula::svSingleVectorRef)
276 : {
277 : const formula::SingleVectorRefToken* pTmpDVR1 =
278 0 : static_cast<const formula::SingleVectorRefToken*>(vSubArguments[i]->GetFormulaToken());
279 0 : ss << " if(singleIndex>=";
280 0 : ss << pTmpDVR1->GetArrayLength();
281 0 : ss << " ||";
282 0 : ss << "isNan(";
283 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(true);
284 0 : ss << "))\n";
285 0 : ss << " tmp";
286 0 : ss << i;
287 0 : ss << "=0;\n else \n";
288 0 : ss << " tmp";
289 0 : ss << i;
290 0 : ss << "=";
291 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(true);
292 0 : ss << ";\n";
293 : }
294 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() ==
295 : formula::svDoubleVectorRef)
296 : {
297 : const formula::DoubleVectorRefToken* pTmpDVR2 =
298 0 : static_cast<const formula::DoubleVectorRefToken*>(vSubArguments[i]->GetFormulaToken());
299 0 : ss << " if(doubleIndex>=";
300 0 : ss << pTmpDVR2->GetArrayLength();
301 0 : ss << " ||";
302 0 : ss << "isNan(";
303 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(false);
304 0 : ss << "))\n";
305 0 : ss << " tmp";
306 0 : ss << i;
307 0 : ss << "=0;\n else \n";
308 0 : ss << " tmp";
309 0 : ss << i;
310 0 : ss << "=";
311 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef(false);
312 0 : ss << ";\n";
313 : }
314 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble ||
315 0 : vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush)
316 : {
317 0 : ss << " if(";
318 0 : ss << "isNan(";
319 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef();
320 0 : ss << "))\n";
321 0 : ss << " tmp";
322 0 : ss << i;
323 0 : ss << "=0;\n else \n";
324 0 : ss << " tmp";
325 0 : ss << i;
326 0 : ss << "=";
327 0 : ss << vSubArguments[i]->GenSlidingWindowDeclRef();
328 0 : ss << ";\n";
329 :
330 : }
331 :
332 0 : }
333 :
334 0 : void CheckVariables::CheckSubArgumentIsNan2( std::stringstream& ss,
335 : SubArguments& vSubArguments, int argumentNum, std::string p )
336 : {
337 0 : int i = argumentNum;
338 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble)
339 : {
340 0 : ss << " tmp";
341 0 : ss << i;
342 0 : ss << "=";
343 0 : vSubArguments[i]->GenDeclRef(ss);
344 0 : ss << ";\n";
345 0 : return;
346 : }
347 :
348 : #ifdef ISNAN
349 0 : ss << " tmp";
350 0 : ss << i;
351 0 : ss << "= fsum(";
352 0 : vSubArguments[i]->GenDeclRef(ss);
353 0 : if (vSubArguments[i]->GetFormulaToken()->GetType() ==
354 : formula::svDoubleVectorRef)
355 0 : ss << "[" << p.c_str() << "]";
356 0 : else if (vSubArguments[i]->GetFormulaToken()->GetType() ==
357 : formula::svSingleVectorRef)
358 0 : ss << "[get_group_id(1)]";
359 0 : ss << ", 0);\n";
360 : #else
361 : ss << " tmp";
362 : ss << i;
363 : ss << "=";
364 : vSubArguments[i]->GenDeclRef(ss);
365 : if (vSubArguments[i]->GetFormulaToken()->GetType() ==
366 : formula::svDoubleVectorRef)
367 : ss << "[" << p.c_str() << "]";
368 : else if (vSubArguments[i]->GetFormulaToken()->GetType() ==
369 : formula::svSingleVectorRef)
370 : ss << "[get_group_id(1)]";
371 :
372 : ss << ";\n";
373 : #endif
374 : }
375 :
376 0 : void CheckVariables::CheckAllSubArgumentIsNan(
377 : std::stringstream& ss, SubArguments& vSubArguments )
378 : {
379 0 : ss << " int k = gid0;\n";
380 0 : for (unsigned i = 0; i < vSubArguments.size(); i++)
381 : {
382 0 : CheckSubArgumentIsNan(ss, vSubArguments, i);
383 : }
384 0 : }
385 :
386 0 : void CheckVariables::UnrollDoubleVector( std::stringstream& ss,
387 : std::stringstream& unrollstr, const formula::DoubleVectorRefToken* pCurDVR,
388 : int nCurWindowSize )
389 : {
390 0 : int unrollSize = 16;
391 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed())
392 : {
393 0 : ss << " loop = (" << nCurWindowSize << " - gid0)/";
394 0 : ss << unrollSize << ";\n";
395 : }
396 0 : else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
397 : {
398 0 : ss << " loop = (" << nCurWindowSize << " + gid0)/";
399 0 : ss << unrollSize << ";\n";
400 :
401 : }
402 : else
403 : {
404 0 : ss << " loop = " << nCurWindowSize << "/" << unrollSize << ";\n";
405 : }
406 :
407 0 : ss << " for ( int j = 0;j< loop; j++)\n";
408 0 : ss << " {\n";
409 0 : ss << " int i = ";
410 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed())
411 : {
412 0 : ss << "gid0 + j * " << unrollSize << ";\n";
413 : }
414 : else
415 : {
416 0 : ss << "j * " << unrollSize << ";\n";
417 : }
418 :
419 0 : if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
420 : {
421 0 : ss << " int doubleIndex = i+gid0;\n";
422 : }
423 : else
424 : {
425 0 : ss << " int doubleIndex = i;\n";
426 : }
427 :
428 0 : for (int j = 0; j < unrollSize; j++)
429 : {
430 0 : ss << unrollstr.str();
431 0 : ss << "i++;\n";
432 0 : ss << "doubleIndex++;\n";
433 : }
434 0 : ss << " }\n";
435 0 : ss << " for (int i = ";
436 0 : if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed())
437 : {
438 0 : ss << "gid0 + loop *" << unrollSize << "; i < ";
439 0 : ss << nCurWindowSize << "; i++)\n";
440 : }
441 0 : else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
442 : {
443 0 : ss << "0 + loop *" << unrollSize << "; i < gid0+";
444 0 : ss << nCurWindowSize << "; i++)\n";
445 : }
446 : else
447 : {
448 0 : ss << "0 + loop *" << unrollSize << "; i < ";
449 0 : ss << nCurWindowSize << "; i++)\n";
450 : }
451 0 : ss << " {\n";
452 0 : if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
453 : {
454 0 : ss << " int doubleIndex = i+gid0;\n";
455 : }
456 : else
457 : {
458 0 : ss << " int doubleIndex = i;\n";
459 : }
460 0 : ss << unrollstr.str();
461 0 : ss << " }\n";
462 0 : }
463 :
464 : }}
465 :
466 : /* vim:set shiftwidth=4 softtabstop=4 expandtab: */
|
