sc/source/core/opencl/opbase.cxx

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