LCOV - code coverage report
Current view: top level - sc/source/core/opencl - opbase.cxx (source / functions) Hit Total Coverage
Test: commit e02a6cb2c3e2b23b203b422e4e0680877f232636 Lines: 0 156 0.0 %
Date: 2014-04-14 Functions: 0 16 0.0 %
Legend: Lines: hit not hit

          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: */

Generated by: LCOV version 1.10