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