sc/source/core/opencl/formulagroupcl.cxx | 224 +++++++++++++++++++++---------- sc/source/core/opencl/op_logical.cxx | 63 ++++++++ sc/source/core/opencl/op_logical.hxx | 7 sc/source/core/opencl/op_math.cxx | 93 ++++-------- sc/source/core/opencl/op_spreadsheet.cxx | 49 +++++- sc/source/core/opencl/op_spreadsheet.hxx | 1 sc/source/core/opencl/op_statistical.cxx | 92 +++++++----- sc/source/core/opencl/opbase.cxx | 54 +++++-- sc/source/core/opencl/opbase.hxx | 9 + 9 files changed, 402 insertions(+), 190 deletions(-)
New commits: commit 4d7e4e63cf4092f0641a26a58fd73dc005d24275 Author: haochen <haoc...@multicorewareinc.com> Date: Fri Jun 13 11:26:12 2014 +0800 GPU Calc:Support nested formulae expansion for simple nested Change-Id: If1ae42a5481cf76942ff1ac5e0ee31a94159badd diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index 5f6b4ff..059927f 100644 --- a/sc/source/core/opencl/formulagroupcl.cxx +++ b/sc/source/core/opencl/formulagroupcl.cxx @@ -1256,12 +1256,12 @@ public: ss << ", "; vSubArguments[i]->GenSlidingWindowDecl(ss); } - ss << ") {\n\t"; - ss << "double tmp = " << GetBottom() <<";\n\t"; - ss << "int gid0 = get_global_id(0);\n\t"; + ss << ") {\n"; + ss << "double tmp = " << GetBottom() <<";\n"; + ss << "int gid0 = get_global_id(0);\n"; if (isAverage()) - ss << "int nCount = 0;\n\t"; - ss << "double tmpBottom;\n\t"; + ss << "int nCount = 0;\n"; + ss << "double tmpBottom;\n"; unsigned i = vSubArguments.size(); while (i--) { @@ -1292,60 +1292,52 @@ public: if (pCur->GetType() == formula::svSingleVectorRef) { -#ifdef ISNAN const formula::SingleVectorRefToken* pSVR = static_cast< const formula::SingleVectorRefToken* >(pCur); - ss << "if (gid0 < " << pSVR->GetArrayLength() << "){\n\t\t"; -#else -#endif + ss << "if (gid0 < " << pSVR->GetArrayLength() << "){\n"; } else if (pCur->GetType() == formula::svDouble) { -#ifdef ISNAN - ss << "{\n\t\t"; -#endif - } - else - { + ss << "{\n"; } } -#ifdef ISNAN if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode()) { - ss << "tmpBottom = " << GetBottom() << ";\n\t\t"; + ss << "tmpBottom = " << GetBottom() << ";\n"; ss << "if (isNan("; ss << vSubArguments[i]->GenSlidingWindowDeclRef(); - ss << "))\n\t\t\t"; - ss << "tmp = "; - ss << Gen2("tmpBottom", "tmp") << ";\n\t\t"; - ss << "else{\n\t\t\t"; - ss << "tmp = "; + ss << "))\n"; + if ( ZeroReturnZero() ) + ss << " return 0;\n"; + else + { + ss << " tmp = "; + ss << Gen2("tmpBottom", "tmp") << ";\n"; + } + ss << "else{\n"; + ss << " tmp = "; ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp"); - ss << ";\n\t\t\t"; - ss << "}\n\t"; - ss << "}\n\t"; + ss << ";\n"; + ss << " }\n"; + ss << "}\n"; + if ( vSubArguments[i]->GetFormulaToken()->GetType() == + formula::svSingleVectorRef&& ZeroReturnZero() ) + { + ss << "else{\n"; + ss << " return 0;\n"; + ss << " }\n"; + } } else { ss << "tmp = "; ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp"); - ss << ";\n\t"; + ss << ";\n"; } -#else - ss << "tmp = "; - // Generate the operation in binary form - ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp"); - ss << ";\n\t"; -#endif } ss << "return tmp"; -#ifdef ISNAN if (isAverage()) ss << "*pow((double)nCount,-1.0)"; -#else - if (isAverage()) - ss << "/(double)"<<nItems; -#endif ss << ";\n}"; } virtual bool isAverage() const { return false; } @@ -1535,6 +1527,11 @@ public: static_cast< const formula::SingleVectorRefToken*> (vSubArguments[i]->GetFormulaToken()); temp3<<pSVR->GetArrayLength(); + temp3 << ")||isNan("<<vSubArguments[i] + ->GenSlidingWindowDeclRef(); + temp3 << ")?0:"; + temp3 << vSubArguments[i]->GenSlidingWindowDeclRef(); + temp3 << ")"; } else if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDoubleVectorRef){ @@ -1542,12 +1539,13 @@ public: static_cast< const formula::DoubleVectorRefToken*> (vSubArguments[i]->GetFormulaToken()); temp3<<pSVR->GetArrayLength(); - } - temp3 << ")||isNan("<<vSubArguments[i] + temp3 << ")||isNan("<<vSubArguments[i] ->GenSlidingWindowDeclRef(true); temp3 << ")?0:"; temp3 << vSubArguments[i]->GenSlidingWindowDeclRef(true); temp3 << ")"; + } + } else temp3 << vSubArguments[i]->GenSlidingWindowDeclRef(true); @@ -1604,6 +1602,11 @@ public: static_cast< const formula::SingleVectorRefToken*> (vSubArguments[i]->GetFormulaToken()); temp4<<pSVR->GetArrayLength(); + temp4 << ")||isNan("<<vSubArguments[i] + ->GenSlidingWindowDeclRef(); + temp4 << ")?0:"; + temp4 << vSubArguments[i]->GenSlidingWindowDeclRef(); + temp4 << ")"; } else if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDoubleVectorRef) @@ -1612,12 +1615,13 @@ public: static_cast< const formula::DoubleVectorRefToken*> (vSubArguments[i]->GetFormulaToken()); temp4<<pSVR->GetArrayLength(); + temp4 << ")||isNan("<<vSubArguments[i] + ->GenSlidingWindowDeclRef(true); + temp4 << ")?0:"; + temp4 << vSubArguments[i]->GenSlidingWindowDeclRef(true); + temp4 << ")"; } - temp4 << ")||isNan("<<vSubArguments[i] - ->GenSlidingWindowDeclRef(true); - temp4 << ")?0:"; - temp4 << vSubArguments[i]->GenSlidingWindowDeclRef(true); - temp4 << ")"; + } else { @@ -1752,6 +1756,7 @@ public: return lhs + "*" + rhs; } virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "fmul"; } + virtual bool ZeroReturnZero() {return true;}; }; /// Technically not a reduction, but fits the framework. @@ -2020,9 +2025,20 @@ public: ss << ")"; } else { if (mvSubArguments.size() != 2) - throw Unhandled(); - ss << "(" << mpCodeGen->Gen2(mvSubArguments[0]->GenSlidingWindowDeclRef(true), - mvSubArguments[1]->GenSlidingWindowDeclRef(true)) << ")"; + throw Unhandled(); + bool bArgument1_NeedNested = + (mvSubArguments[0]->GetFormulaToken()->GetType() + == formula::svSingleVectorRef)? false:true; + bool bArgument2_NeedNested = + (mvSubArguments[1]->GetFormulaToken()->GetType() + == formula::svSingleVectorRef) ? false:true; + ss << "("; + ss << mpCodeGen-> + Gen2(mvSubArguments[0] + ->GenSlidingWindowDeclRef(bArgument1_NeedNested), + mvSubArguments[1] + ->GenSlidingWindowDeclRef(bArgument2_NeedNested)); + ss << ")"; } return ss.str(); } diff --git a/sc/source/core/opencl/op_logical.cxx b/sc/source/core/opencl/op_logical.cxx index 9dce77b..96170f2 100644 --- a/sc/source/core/opencl/op_logical.cxx +++ b/sc/source/core/opencl/op_logical.cxx @@ -319,7 +319,6 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss, { ss << "\ndouble " << sSymName; ss << "_"<< BinFuncName() <<"("; - if(vSubArguments.size()!=3) throw Unhandled("unknown operand for ocPush"); for (unsigned i = 0; i < vSubArguments.size(); i++) { if (i) @@ -332,22 +331,49 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss, FormulaToken *tmpCur0 = vSubArguments[0]->GetFormulaToken(); if(tmpCur0->GetType() == formula::svDoubleVectorRef) { - throw UnhandledToken(tmpCur0, "unknown operand for ocPush"); + throw UnhandledToken(tmpCur0, "unknown operand for ocPush"); } else { - ss << " if(isNan("; - ss << vSubArguments[0]->GenSlidingWindowDeclRef(); - ss << ")|| "; - ss << vSubArguments[0]->GenSlidingWindowDeclRef(); - ss << " == 0)\n"; - ss << " return "; - ss << vSubArguments[2]->GenSlidingWindowDeclRef(); - ss << ";\n"; - ss << " else"; - ss <<" return "; - ss << vSubArguments[1]->GenSlidingWindowDeclRef(); - ss <<";\n"; + if(vSubArguments.size()==3) + { + ss << " if(isNan("; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << ")|| "; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << " == 0)\n"; + ss << " return "; + ss << vSubArguments[2]->GenSlidingWindowDeclRef(); + ss << ";\n"; + ss << " else"; + ss <<" return "; + ss << vSubArguments[1]->GenSlidingWindowDeclRef(); + ss <<";\n"; + } + if(vSubArguments.size()==2) + { + ss << " if(isNan("; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << ")|| "; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << " == 0)\n"; + ss << " return 0;\n"; + ss << " else"; + ss <<" return "; + ss << vSubArguments[1]->GenSlidingWindowDeclRef(); + ss <<";\n"; + } + if(vSubArguments.size()==1) + { + ss << " if(isNan("; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << ")|| "; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << " == 0)\n"; + ss << " return 0;\n"; + ss << " else"; + ss <<" return 1;\n"; + } } ss << "}\n"; } diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx index 02b9c37..b0751f5 100644 --- a/sc/source/core/opencl/opbase.cxx +++ b/sc/source/core/opencl/opbase.cxx @@ -130,7 +130,6 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, SubArguments &vSubArguments, int argumentNum) { int i = argumentNum; -#ifdef ISNAN if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svSingleVectorRef) { @@ -139,6 +138,17 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ss<< " if(singleIndex>="; ss<< pTmpDVR1->GetArrayLength(); ss<<" ||"; + ss<< "isNan("; + ss<< vSubArguments[i]->GenSlidingWindowDeclRef(true); + ss<<"))\n"; + ss<< " tmp"; + ss<< i; + ss <<"=0;\n else \n"; + ss <<" tmp"; + ss <<i; + ss << "="; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(true); + ss<<";\n"; } if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDoubleVectorRef) @@ -148,24 +158,36 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ss<< " if(doubleIndex>="; ss<< pTmpDVR2->GetArrayLength(); ss<<" ||"; + ss<< "isNan("; + ss<< vSubArguments[i]->GenSlidingWindowDeclRef(false); + ss<<"))\n"; + ss<< " tmp"; + ss<< i; + ss <<"=0;\n else \n"; + ss <<" tmp"; + ss <<i; + ss << "="; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(false); + ss<<";\n"; } if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble || vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush) { ss<< " if("; + ss<< "isNan("; + ss<< vSubArguments[i]->GenSlidingWindowDeclRef(); + ss<<"))\n"; + ss<< " tmp"; + ss<< i; + ss <<"=0;\n else \n"; + ss <<" tmp"; + ss <<i; + ss << "="; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(); + ss<<";\n"; + } - ss<< "isNan("; - ss<< vSubArguments[i]->GenSlidingWindowDeclRef(true); - ss<<"))\n"; - ss<< " tmp"; - ss<< i; - ss <<"=0;\n else \n"; -#endif - ss <<" tmp"; - ss <<i; - ss << "="; - ss << vSubArguments[i]->GenSlidingWindowDeclRef(true); - ss<<";\n"; + } void CheckVariables::CheckSubArgumentIsNan2( std::stringstream & ss, diff --git a/sc/source/core/opencl/opbase.hxx b/sc/source/core/opencl/opbase.hxx index d1d0005..487fc04 100644 --- a/sc/source/core/opencl/opbase.hxx +++ b/sc/source/core/opencl/opbase.hxx @@ -211,7 +211,7 @@ public: virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE; /// When referenced in a sliding window function - virtual std::string GenSlidingWindowDeclRef(bool=true) const SAL_OVERRIDE; + virtual std::string GenSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE; /// Create buffer and pass the buffer to a given kernel virtual size_t Marshal(cl_kernel, int, int, cl_program) SAL_OVERRIDE; @@ -250,6 +250,9 @@ public: std::set<std::string>& ) {} virtual bool takeString() const = 0; virtual bool takeNumeric() const = 0; + //Continue process 'Zero' or Not(like OpMul, not continue process when meet + // 'Zero' + virtual bool ZeroReturnZero() {return false;} virtual ~OpBase() {} }; commit c1dfbcdd114929cb5f03972e419e0452c3bbd9ef Author: haochen <haoc...@multicorewareinc.com> Date: Wed Jun 11 12:59:37 2014 +0800 GPU Calc:Add more judge for NAN in SingleVector parameter Change-Id: I01f2576e9d8f6a2e677d1fb08097bc81f15bfbe0 diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx index ccd653c..02b9c37 100644 --- a/sc/source/core/opencl/opbase.cxx +++ b/sc/source/core/opencl/opbase.cxx @@ -64,10 +64,10 @@ std::string VectorRef::GenSlidingWindowDeclRef(bool nested) const std::stringstream ss; formula::SingleVectorRefToken *pSVR = dynamic_cast<formula::SingleVectorRefToken*>(DynamicKernelArgument::GetFormulaToken()); - if (pSVR&&nested) + if (pSVR&&!nested) ss << "(gid0 < " << pSVR->GetArrayLength() << "?"; ss << mSymName << "[gid0]"; - if (pSVR&&nested) + if (pSVR&&!nested) ss << ":NAN)"; return ss.str(); } @@ -155,7 +155,7 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ss<< " if("; } ss<< "isNan("; - ss<< vSubArguments[i]->GenSlidingWindowDeclRef(false); + ss<< vSubArguments[i]->GenSlidingWindowDeclRef(true); ss<<"))\n"; ss<< " tmp"; ss<< i; @@ -164,7 +164,7 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ss <<" tmp"; ss <<i; ss << "="; - ss << vSubArguments[i]->GenSlidingWindowDeclRef(false); + ss << vSubArguments[i]->GenSlidingWindowDeclRef(true); ss<<";\n"; } commit 848c76d5e3e9b477dbce735d8a0bd734cd5e389c Author: haochen <haoc...@multicorewarein.com> Date: Sat May 31 14:28:22 2014 +0800 GPU Calc:Support default 2nd parameter in ROUND Change-Id: I003ba9c945dbc3c6417d0502902610c0eaff2bda diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index 9b45c91..5f6b4ff 100644 --- a/sc/source/core/opencl/formulagroupcl.cxx +++ b/sc/source/core/opencl/formulagroupcl.cxx @@ -1487,7 +1487,9 @@ public: #ifdef UNROLLING_FACTOR ss << "\tint i;\n\t"; - ss << "int currentCount0, currentCount1;\n\t"; + ss << "int currentCount0;\n"; + for ( unsigned i = 0; i < vSubArguments.size()-1; i++) + ss << "int currentCount"<<i+1<<";\n"; std::stringstream temp3,temp4; int outLoopSize = UNROLLING_FACTOR; if (nCurWindowSize/outLoopSize != 0){ diff --git a/sc/source/core/opencl/op_math.cxx b/sc/source/core/opencl/op_math.cxx index 36c94f6..555da592 100644 --- a/sc/source/core/opencl/op_math.cxx +++ b/sc/source/core/opencl/op_math.cxx @@ -1707,11 +1707,17 @@ void OpRound::GenSlidingWindowFunction(std::stringstream &ss, ss << " int singleIndex = gid0;\n"; GenTmpVariables(ss,vSubArguments); CheckAllSubArgumentIsNan(ss,vSubArguments); - ss << " for(int i=0;i<tmp1;i++)\n"; - ss << " tmp0 = tmp0 * 10;\n"; + if(vSubArguments.size() ==2) + { + ss << " for(int i=0;i<tmp1;i++)\n"; + ss << " tmp0 = tmp0 * 10;\n"; + } ss << " double tmp=round(tmp0);\n"; - ss << " for(int i=0;i<tmp1;i++)\n"; - ss << " tmp = tmp / 10;\n"; + if(vSubArguments.size() ==2) + { + ss << " for(int i=0;i<tmp1;i++)\n"; + ss << " tmp = tmp / 10;\n"; + } ss << " return tmp;\n"; ss << "}"; } commit 47a10f026622849ae22d1f559329520ba9fc197d Author: haochen <haoc...@multicorewareinc.com> Date: Sat May 31 13:54:22 2014 +0800 GPU Calc:Support 3rd parameter in FLOOR Change-Id: Ie3a265f34a5f589d41e802b63df4be6a64989b05 diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index 91a54f4..9b45c91 100644 --- a/sc/source/core/opencl/formulagroupcl.cxx +++ b/sc/source/core/opencl/formulagroupcl.cxx @@ -419,7 +419,7 @@ size_t DynamicKernelStringArgument::Marshal(cl_kernel k, int argno, int, cl_prog if (CL_SUCCESS != err) throw OpenCLError(err, __FILE__, __LINE__); - cl_uint *pHashBuffer = (cl_uint*)clEnqueueMapBuffer( + pHashBuffer = (cl_uint*)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0, szHostBuffer, 0, NULL, NULL, &err); if (CL_SUCCESS != err) diff --git a/sc/source/core/opencl/op_math.cxx b/sc/source/core/opencl/op_math.cxx index 4a960b3..36c94f6 100644 --- a/sc/source/core/opencl/op_math.cxx +++ b/sc/source/core/opencl/op_math.cxx @@ -424,14 +424,6 @@ void OpEven::GenSlidingWindowFunction(std::stringstream &ss, void OpMod::GenSlidingWindowFunction(std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments) { -#ifdef ISNAN - FormulaToken *tmpCur0 = vSubArguments[0]->GetFormulaToken(); - const formula::SingleVectorRefToken*tmpCurDVR0= static_cast<const - formula::SingleVectorRefToken *>(tmpCur0); - FormulaToken *tmpCur1 = vSubArguments[0]->GetFormulaToken(); - const formula::SingleVectorRefToken*tmpCurDVR1= static_cast<const - formula::SingleVectorRefToken *>(tmpCur1); -#endif ss << "\ndouble " << sSymName; ss << "_"<< BinFuncName() <<"("; for (unsigned i = 0; i < vSubArguments.size(); i++) @@ -444,22 +436,13 @@ void OpMod::GenSlidingWindowFunction(std::stringstream &ss, ss <<" int gid0=get_global_id(0);\n"; ss << " double arg0 = " << vSubArguments[0]->GenSlidingWindowDeclRef(); ss << ";\n"; - ss <<" double arg1 =" << vSubArguments[1]->GenSlidingWindowDeclRef(); + ss << " double arg1 =" << vSubArguments[1]->GenSlidingWindowDeclRef(); ss << ";\n"; -#ifdef ISNAN - ss<< " if(isNan(arg0)||(gid0>="; - ss<<tmpCurDVR0->GetArrayLength(); - ss<<"))\n"; - ss<<" arg0 = 0;\n"; -#endif -#ifdef ISNAN - ss<< " if(isNan(arg1)||(gid0>="; - ss<<tmpCurDVR1->GetArrayLength(); - ss<<"))\n"; - ss<<" arg1 = 0;\n"; -#endif + ss << " if(isNan(arg0)||arg0 == 0)\n"; + ss << " return 0;\n"; + ss << " if(isNan(arg1) || arg1 ==0)\n"; + ss << " return NAN;\n"; ss << " double tem;\n"; - ss << " if(arg1 != 0) {\n"; ss << " if(arg0 < 0 && arg1 > 0)\n"; ss << " while(arg0 < 0)\n"; ss << " arg0 += arg1;\n"; @@ -467,9 +450,6 @@ void OpMod::GenSlidingWindowFunction(std::stringstream &ss, ss << " while(arg0 > 0)\n"; ss << " arg0 += arg1;\n"; ss << " tem = fmod(arg0,arg1);\n"; - ss << " }\n"; - ss << " else\n"; - ss << " tem = 0;\n"; ss << " if(arg1 < 0 && tem > 0)\n"; ss << " tem = -tem;\n"; ss << " return tem;\n"; @@ -2318,9 +2298,6 @@ void OpFloor::GenSlidingWindowFunction( std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments) { - FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken(); - const formula::SingleVectorRefToken*tmpCurDVR= static_cast<const - formula::SingleVectorRefToken *>(tmpCur); ss << "\ndouble " << sSymName; ss << "_"<< BinFuncName() <<"("; for (unsigned i = 0; i < vSubArguments.size(); i++) @@ -2330,35 +2307,27 @@ void OpFloor::GenSlidingWindowFunction( vSubArguments[i]->GenSlidingWindowDecl(ss); } ss << ")\n{\n"; - ss <<" int gid0=get_global_id(0);\n"; - ss << " double arg0 = " << vSubArguments[0]->GenSlidingWindowDeclRef(); - ss << ";\n"; - ss << " double arg1 = " << vSubArguments[1]->GenSlidingWindowDeclRef(); + ss << " int gid0=get_global_id(0);\n"; + ss << " double arg0,arg1,arg2=0.0;\n"; + ss << " arg0 = " << vSubArguments[0]->GenSlidingWindowDeclRef(); ss << ";\n"; - ss << " double arg2 = " << vSubArguments[2]->GenSlidingWindowDeclRef(); + ss << " arg1 = " << vSubArguments[1]->GenSlidingWindowDeclRef(); ss << ";\n"; -#ifdef ISNAN - ss<< " if(isNan(arg0)||(gid0>="; - ss<<tmpCurDVR->GetArrayLength(); - ss<<"))\n"; - ss<<" arg0 = 0;\n"; - ss<< " if(isNan(arg1)||(gid0>="; - ss<<tmpCurDVR->GetArrayLength(); - ss<<"))\n"; - ss<<" arg1 = 0;\n"; - ss<< " if(isNan(arg2)||(gid0>="; - ss<<tmpCurDVR->GetArrayLength(); - ss<<"))\n"; - ss<<" arg2 = 0;\n"; -#endif - ss <<" if(arg1==0.0)\n"; - ss <<" return 0.0;\n"; - ss <<" else if(arg0*arg1<0.0)\n"; - ss <<" return 0.0000000001;\n"; - ss <<" else if(arg2==0.0&&arg0<0.0)\n"; - ss <<" return (trunc(arg0/arg1)+1)*arg1;\n"; - ss <<" else\n"; - ss <<" return trunc(arg0/arg1)*arg1;\n"; + if ( 3 == vSubArguments.size() ) + { + ss << " arg2 = " << vSubArguments[2]->GenSlidingWindowDeclRef(); + ss << ";\n"; + } + ss << " if(isNan(arg0) || isNan(arg1))\n"; + ss << " return 0;\n"; + ss << " if(isNan(arg2))\n"; + ss << " arg2 = 0.0;\n"; + ss << " if(arg0*arg1<0)\n"; + ss << " return NAN;\n"; + ss << " else if(arg2==0.0&&arg0<0.0)\n"; + ss << " return (trunc(arg0/arg1)+1)*arg1;\n"; + ss << " else\n"; + ss << " return trunc(arg0/arg1)*arg1;\n"; ss << "}\n"; } void OpBitOr::GenSlidingWindowFunction(std::stringstream &ss, diff --git a/sc/source/core/opencl/op_statistical.cxx b/sc/source/core/opencl/op_statistical.cxx index 88541e2..20dbb49 100644 --- a/sc/source/core/opencl/op_statistical.cxx +++ b/sc/source/core/opencl/op_statistical.cxx @@ -3678,6 +3678,20 @@ void OpGamma::GenSlidingWindowFunction( void OpCorrel::GenSlidingWindowFunction( std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments) { + if( vSubArguments.size() !=2 ||vSubArguments[0]->GetFormulaToken() + ->GetType() != formula::svDoubleVectorRef||vSubArguments[1] + ->GetFormulaToken()->GetType() != formula::svDoubleVectorRef ) + ///only support DoubleVector in OpCorrelfor GPU calculating. + throw Unhandled(); + const formula::DoubleVectorRefToken* pCurDVRX = + static_cast<const formula::DoubleVectorRefToken *>( + vSubArguments[0]->GetFormulaToken()); + const formula::DoubleVectorRefToken* pCurDVRY = + static_cast<const formula::DoubleVectorRefToken *>( + vSubArguments[1]->GetFormulaToken()); + if( pCurDVRX->GetRefRowSize() != pCurDVRY->GetRefRowSize() ) + throw Unhandled(); + ss << "\ndouble " << sSymName; ss << "_"<< BinFuncName() <<"("; for (unsigned i = 0; i < vSubArguments.size(); i++) @@ -3698,16 +3712,8 @@ void OpCorrel::GenSlidingWindowFunction( ss << "double arg1 = 0.0;\n\t"; ss << "int cnt = 0;\n\t"; - FormulaToken *pCurX = vSubArguments[0]->GetFormulaToken(); - FormulaToken *pCurY = vSubArguments[1]->GetFormulaToken(); - const formula::DoubleVectorRefToken* pCurDVRX = - static_cast<const formula::DoubleVectorRefToken *>(pCurX); - const formula::DoubleVectorRefToken* pCurDVRY = - static_cast<const formula::DoubleVectorRefToken *>(pCurY); - size_t nCurWindowSizeX = pCurDVRX->GetRefRowSize(); - size_t nCurWindowSizeY = pCurDVRY->GetRefRowSize(); - if(nCurWindowSizeX == nCurWindowSizeY) - { + size_t nCurWindowSizeX = pCurDVRY->GetRefRowSize(); + ss << "for (int i = "; if (!pCurDVRX->IsStartFixed() && pCurDVRX->IsEndFixed()) { ss << "gid0; i < " << nCurWindowSizeX << "; i++) {\n\t\t"; @@ -3880,7 +3886,6 @@ void OpCorrel::GenSlidingWindowFunction( ss << "}\n\t"; ss << "}\n"; ss << "}"; - } } void OpNegbinomdist::GenSlidingWindowFunction( @@ -3959,10 +3964,20 @@ void OpNegbinomdist::GenSlidingWindowFunction( void OpPearson::GenSlidingWindowFunction( std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments) { - FormulaToken* pCur = vSubArguments[0]->GetFormulaToken(); - assert(pCur); + if( vSubArguments.size() !=2 ||vSubArguments[0]->GetFormulaToken() + ->GetType() != formula::svDoubleVectorRef||vSubArguments[1] + ->GetFormulaToken()->GetType() != formula::svDoubleVectorRef ) + ///only support DoubleVector in OpPearson for GPU calculating. + throw Unhandled(); const formula::DoubleVectorRefToken* pDVR = - static_cast<const formula::DoubleVectorRefToken *>(pCur); + static_cast<const formula::DoubleVectorRefToken *>( + vSubArguments[0]->GetFormulaToken()); + const formula::DoubleVectorRefToken* pCurDVRY = + static_cast<const formula::DoubleVectorRefToken *>( + vSubArguments[1]->GetFormulaToken()); + if( pDVR->GetRefRowSize() != pCurDVRY->GetRefRowSize() ) + throw Unhandled(); + size_t nCurWindowSize = pDVR->GetRefRowSize(); ss << "\ndouble " << sSymName; @@ -4000,6 +4015,7 @@ void OpPearson::GenSlidingWindowFunction( ss << ";\n"; ss << " fIny = "<<vSubArguments[1]->GenSlidingWindowDeclRef(true); ss << " ;\n"; + ss << " if(isNan(fInx)||isNan(fIny)){fInx=0.0;fIny=0.0;fCount = fCount-1;}\n"; ss << " fSumX += fInx;\n"; ss << " fSumY += fIny;\n"; ss << " fCount = fCount + 1;\n"; @@ -4026,6 +4042,7 @@ void OpPearson::GenSlidingWindowFunction( ss << " ;\n"; ss << " fIny = "<<vSubArguments[1]->GenSlidingWindowDeclRef(true); ss << " ;\n"; + ss << " if(isNan(fInx)||isNan(fIny)){fInx=0.0;fIny=0.0;}\n"; ss << " fSumDeltaXDeltaY += (fInx - fMeanX) * (fIny - fMeanY);\n"; ss << " fSumX += pow(fInx - fMeanX,2);\n"; ss << " fSumY += pow(fIny - fMeanY,2);\n"; @@ -4579,11 +4596,21 @@ void OpCritBinom::GenSlidingWindowFunction(std::stringstream& ss, void OpRsq::GenSlidingWindowFunction( std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments) { - FormulaToken* pCur = vSubArguments[1]->GetFormulaToken(); - assert(pCur); - const formula::DoubleVectorRefToken* pCurDVR = - static_cast<const formula::DoubleVectorRefToken *>(pCur); - size_t nCurWindowSize = pCurDVR->GetRefRowSize(); + if( vSubArguments.size() !=2 ||vSubArguments[0]->GetFormulaToken() + ->GetType() != formula::svDoubleVectorRef||vSubArguments[1] + ->GetFormulaToken()->GetType() != formula::svDoubleVectorRef ) + ///only support DoubleVector in OpRsq for GPU calculating. + throw Unhandled(); + const formula::DoubleVectorRefToken* pCurDVR1 = + static_cast<const formula::DoubleVectorRefToken *>( + vSubArguments[0]->GetFormulaToken()); + const formula::DoubleVectorRefToken* pCurDVR2 = + static_cast<const formula::DoubleVectorRefToken *>( + vSubArguments[1]->GetFormulaToken()); + if( pCurDVR1->GetRefRowSize() != pCurDVR2->GetRefRowSize() ) + throw Unhandled(); + + size_t nCurWindowSize = pCurDVR1->GetRefRowSize(); ss << "\ndouble " << sSymName; ss << "_"<< BinFuncName() <<"("; @@ -4605,29 +4632,18 @@ void OpRsq::GenSlidingWindowFunction( ss << " double tmp0,tmp1;\n"; vSubArguments.size(); ss <<"\n"; - FormulaToken *tmpCur0 = vSubArguments[0]->GetFormulaToken(); - const formula::DoubleVectorRefToken*tmpCurDVR0= static_cast<const - formula::DoubleVectorRefToken *>(tmpCur0); - FormulaToken *tmpCur1 = vSubArguments[1]->GetFormulaToken(); - const formula::DoubleVectorRefToken*tmpCurDVR1= static_cast<const - formula::DoubleVectorRefToken *>(tmpCur1); - ss << " int buffer_fInx_len = "; - ss << tmpCurDVR0->GetArrayLength(); - ss << ";\n"; - ss << " int buffer_fIny_len = "; - ss << tmpCurDVR1->GetArrayLength(); - ss << ";\n"; + ss << " for(int i=0; i<"<<nCurWindowSize<<"; i++)\n"; ss << " {\n"; - ss << " if((gid0+i)>=buffer_fInx_len || isNan("; - ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << " if(isNan("; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(true); ss << "))\n"; ss << " fInx = 0;\n"; ss << " else\n"; ss << " fInx = "<<vSubArguments[0]->GenSlidingWindowDeclRef(); ss << ";\n"; - ss << " if((gid0+i)>=buffer_fIny_len || isNan("; - ss << vSubArguments[1]->GenSlidingWindowDeclRef(); + ss << " if(isNan("; + ss << vSubArguments[1]->GenSlidingWindowDeclRef(true); ss << "))\n"; ss << " fIny = 0;\n"; ss << " else\n"; @@ -4643,14 +4659,14 @@ void OpRsq::GenSlidingWindowFunction( ss << " fSumY = 0.0;\n"; ss << " for(int i=0; i<"<<nCurWindowSize<<"; i++)\n"; ss << " {\n"; - ss << " if((gid0+i)>=buffer_fInx_len || isNan("; - ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << " if(isNan("; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(true); ss << "))\n"; ss << " fInx = 0;\n"; ss << " else\n"; ss << " fInx = "<<vSubArguments[0]->GenSlidingWindowDeclRef(); ss << ";\n"; - ss << " if((gid0+i)>=buffer_fIny_len || isNan("; + ss << " if(isNan("; ss << vSubArguments[1]->GenSlidingWindowDeclRef(); ss << "))\n"; ss << " fIny = 0;\n"; commit a79817c8b1eff0a9e2ad852601776fe34e76d17a Author: haochen <haoc...@multicorewareinc.com> Date: Thu May 29 09:14:53 2014 +0800 GPU Calc:Support IF formula in GPUInterpret Change-Id: I9b2cebb99812d28e25c961129f73585d60690846 diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index c0b62f2..91a54f4 100644 --- a/sc/source/core/opencl/formulagroupcl.cxx +++ b/sc/source/core/opencl/formulagroupcl.cxx @@ -1682,6 +1682,18 @@ public: } virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "leq"; } }; +class OpLess: public Binary { +public: + virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; } + virtual std::string Gen2(const std::string &lhs, const std::string &rhs) const SAL_OVERRIDE + { + std::stringstream ss; + ss << "("<< lhs << "<" << rhs <<")"; + return ss.str(); + } + virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "less"; } +}; + class OpGreater: public Binary { public: @@ -2267,6 +2279,9 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments( case ocLessEqual: mvSubArguments.push_back(SoPHelper(ts, ft->Children[i], new OpLessEqual)); break; + case ocLess: + mvSubArguments.push_back(SoPHelper(ts, ft->Children[i], new OpLess)); + break; case ocEqual: mvSubArguments.push_back(SoPHelper(ts, ft->Children[i], new OpEqual)); break; @@ -2872,6 +2887,10 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments( mvSubArguments.push_back(SoPHelper(ts, ft->Children[i], new OpAveDev)); break; + case ocIf: + mvSubArguments.push_back(SoPHelper(ts, + ft->Children[i], new OpIf)); + break; case ocExternal: if ( !(pChild->GetExternal().compareTo(OUString( "com.sun.star.sheet.addin.Analysis.getEffect")))) diff --git a/sc/source/core/opencl/op_logical.cxx b/sc/source/core/opencl/op_logical.cxx index 79dc74a..9dce77b 100644 --- a/sc/source/core/opencl/op_logical.cxx +++ b/sc/source/core/opencl/op_logical.cxx @@ -314,6 +314,43 @@ void OpXor::GenSlidingWindowFunction(std::stringstream &ss, ss << " return t;\n"; ss << "}\n"; } +void OpIf::GenSlidingWindowFunction(std::stringstream &ss, + const std::string &sSymName, SubArguments &vSubArguments) +{ + ss << "\ndouble " << sSymName; + ss << "_"<< BinFuncName() <<"("; + if(vSubArguments.size()!=3) throw Unhandled("unknown operand for ocPush"); + for (unsigned i = 0; i < vSubArguments.size(); i++) + { + if (i) + ss << ","; + vSubArguments[i]->GenSlidingWindowDecl(ss); + } + ss << ") {\n"; + ss << " int gid0 = get_global_id(0);\n"; + + FormulaToken *tmpCur0 = vSubArguments[0]->GetFormulaToken(); + if(tmpCur0->GetType() == formula::svDoubleVectorRef) + { + throw UnhandledToken(tmpCur0, "unknown operand for ocPush"); + } + else + { + ss << " if(isNan("; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << ")|| "; + ss << vSubArguments[0]->GenSlidingWindowDeclRef(); + ss << " == 0)\n"; + ss << " return "; + ss << vSubArguments[2]->GenSlidingWindowDeclRef(); + ss << ";\n"; + ss << " else"; + ss <<" return "; + ss << vSubArguments[1]->GenSlidingWindowDeclRef(); + ss <<";\n"; + } + ss << "}\n"; +} }} /* vim:set shiftwidth=4 softtabstop=4 expandtab: */ diff --git a/sc/source/core/opencl/op_logical.hxx b/sc/source/core/opencl/op_logical.hxx index 357ddd8..4365a57 100644 --- a/sc/source/core/opencl/op_logical.hxx +++ b/sc/source/core/opencl/op_logical.hxx @@ -44,6 +44,13 @@ public: const std::string &sSymName, SubArguments &vSubArguments) SAL_OVERRIDE; virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "Xor"; } }; +class OpIf:public Normal +{ +public: + virtual void GenSlidingWindowFunction(std::stringstream &ss, + const std::string &sSymName, SubArguments &vSubArguments) SAL_OVERRIDE; + virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "IF"; } +}; }} commit bf18bb2aababe3425e5a3246e349abdfb94f146f Author: haochen <haoc...@multicorewareinc.com> Date: Thu May 29 09:09:55 2014 +0800 GPU Calc:Support string arguments in VLookup Change-Id: Ic2400a13c07c5b08beccaeffef4899c8f8b43af8 diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index 4f40e25..c0b62f2 100644 --- a/sc/source/core/opencl/formulagroupcl.cxx +++ b/sc/source/core/opencl/formulagroupcl.cxx @@ -378,28 +378,55 @@ size_t DynamicKernelStringArgument::Marshal(cl_kernel k, int argno, int, cl_prog vRef = pDVR->GetArrays()[mnIndex]; } size_t szHostBuffer = nStrings * sizeof(cl_int); - // Marshal strings. Right now we pass hashes of these string - mpClmem = clCreateBuffer(kEnv.mpkContext, - (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, - szHostBuffer, NULL, &err); - if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); - cl_uint *pHashBuffer = (cl_uint*)clEnqueueMapBuffer( - kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0, - szHostBuffer, 0, NULL, NULL, &err); - if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); - for (size_t i = 0; i < nStrings; i++) + cl_uint *pHashBuffer = NULL; + + if ( vRef.mpStringArray != NULL) { - if (vRef.mpStringArray[i]) - { - const OUString tmp = OUString(vRef.mpStringArray[i]); - pHashBuffer[i] = tmp.hashCode(); - } - else + // Marshal strings. Right now we pass hashes of these string + mpClmem = clCreateBuffer(kEnv.mpkContext, + (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, + szHostBuffer, NULL, &err); + if (CL_SUCCESS != err) + throw OpenCLError(err, __FILE__, __LINE__); + + pHashBuffer = (cl_uint*)clEnqueueMapBuffer( + kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0, + szHostBuffer, 0, NULL, NULL, &err); + if (CL_SUCCESS != err) + throw OpenCLError(err, __FILE__, __LINE__); + + for (size_t i = 0; i < nStrings; i++) { + if (vRef.mpStringArray[i]) + { + const OUString tmp = OUString(vRef.mpStringArray[i]); + pHashBuffer[i] = tmp.hashCode(); + } + else + { + pHashBuffer[i] = 0; + } + } + } + else + { + if (nStrings == 0) + szHostBuffer = sizeof(cl_int); // a dummy small value + // Marshal as a buffer of NANs + mpClmem = clCreateBuffer(kEnv.mpkContext, + (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, + szHostBuffer, NULL, &err); + if (CL_SUCCESS != err) + throw OpenCLError(err, __FILE__, __LINE__); + + cl_uint *pHashBuffer = (cl_uint*)clEnqueueMapBuffer( + kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0, + szHostBuffer, 0, NULL, NULL, &err); + if (CL_SUCCESS != err) + throw OpenCLError(err, __FILE__, __LINE__); + + for (size_t i = 0; i < szHostBuffer/sizeof(cl_int); i++) pHashBuffer[i] = 0; - } } err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem, pHashBuffer, 0, NULL, NULL); @@ -425,6 +452,7 @@ public: ss << ", "; mStringArgument.GenSlidingWindowDecl(ss); } + virtual bool IsMixedArgument() const SAL_OVERRIDE {return true;} virtual void GenSlidingWindowFunction(std::stringstream &) SAL_OVERRIDE {} /// Generate declaration virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE @@ -439,12 +467,20 @@ public: ss << ","; mStringArgument.GenDeclRef(ss); } - virtual std::string GenSlidingWindowDeclRef(bool) const SAL_OVERRIDE + virtual void GenNumDeclRef(std::stringstream& ss) const SAL_OVERRIDE + { + VectorRef::GenSlidingWindowDecl(ss); + } + virtual void GenStringDeclRef(std::stringstream& ss) const SAL_OVERRIDE + { + mStringArgument.GenSlidingWindowDecl(ss); + } + virtual std::string GenSlidingWindowDeclRef(bool nested) const SAL_OVERRIDE { std::stringstream ss; ss << "(!isNan(" << VectorRef::GenSlidingWindowDeclRef(); ss << ")?" << VectorRef::GenSlidingWindowDeclRef(); - ss << ":" << mStringArgument.GenSlidingWindowDeclRef(); + ss << ":" << mStringArgument.GenSlidingWindowDeclRef(nested); ss << ")"; return ss.str(); } @@ -725,15 +761,16 @@ public: ss << ","; mStringArgument.GenDeclRef(ss); } - virtual std::string GenSlidingWindowDeclRef(bool) const SAL_OVERRIDE + virtual std::string GenSlidingWindowDeclRef(bool nested) const SAL_OVERRIDE { std::stringstream ss; ss << "(!isNan(" << mDoubleArgument.GenSlidingWindowDeclRef(); ss << ")?" << mDoubleArgument.GenSlidingWindowDeclRef(); - ss << ":" << mStringArgument.GenSlidingWindowDeclRef(); + ss << ":" << mStringArgument.GenSlidingWindowDeclRef(nested); ss << ")"; return ss.str(); } + virtual bool IsMixedArgument() const SAL_OVERRIDE {return true;} virtual std::string GenDoubleSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE { std::stringstream ss; @@ -746,6 +783,14 @@ public: ss << mStringArgument.GenSlidingWindowDeclRef(); return ss.str(); } + virtual void GenNumDeclRef(std::stringstream& ss) const SAL_OVERRIDE + { + mDoubleArgument.GenDeclRef(ss); + } + virtual void GenStringDeclRef(std::stringstream& ss) const SAL_OVERRIDE + { + mStringArgument.GenDeclRef(ss); + } virtual size_t Marshal(cl_kernel k, int argno, int vw, cl_program p) SAL_OVERRIDE { int i = mDoubleArgument.Marshal(k, argno, vw, p); diff --git a/sc/source/core/opencl/op_spreadsheet.cxx b/sc/source/core/opencl/op_spreadsheet.cxx index d2b78c9..3610850 100644 --- a/sc/source/core/opencl/op_spreadsheet.cxx +++ b/sc/source/core/opencl/op_spreadsheet.cxx @@ -157,10 +157,32 @@ void OpVLookup::GenSlidingWindowFunction(std::stringstream &ss, ss << " == "; ss << j+1; ss << ")\n"; - ss << " tmp = "; - vSubArguments[1+j]->GenDeclRef(ss); - ss << "[rowNum];\n"; + if( !(vSubArguments[1+j]->IsMixedArgument())) + { + ss << "{"; + ss << " tmp = "; + vSubArguments[1+j]->GenDeclRef(ss); + ss << "[rowNum];\n"; + ss << "}"; + + } + else + { + ss << "{"; + + ss << " tmp = isNan("; + vSubArguments[1+j]->GenNumDeclRef(ss); + ss << "[rowNum]"<<")?"; + vSubArguments[1+j]->GenNumDeclRef(ss); + ss << "[rowNum]:"; + vSubArguments[1+j]->GenStringDeclRef(ss); + ss << "[rowNum];\n"; + ss << "}"; + + + + } } ss << " return tmp;\n"; ss << " }\n"; @@ -221,9 +243,24 @@ void OpVLookup::GenSlidingWindowFunction(std::stringstream &ss, ss << " == "; ss << j+1; ss << ")\n"; - ss << " tmp = "; - vSubArguments[1+j]->GenDeclRef(ss); - ss << "[rowNum];\n"; + ///Add MixedArguments for string support in Vlookup. + if( !(vSubArguments[1+j]->IsMixedArgument())) + { + ss << " tmp = "; + vSubArguments[1+j]->GenDeclRef(ss); + ss << "[rowNum];\n"; + } + else + { + ss << " tmp = isNan("; + vSubArguments[1+j]->GenNumDeclRef(ss); + ss << "[rowNum]"<<")?"; + vSubArguments[1+j]->GenNumDeclRef(ss); + ss << "[rowNum]:"; + vSubArguments[1+j]->GenStringDeclRef(ss); + ss << "[rowNum];\n"; + } + } ss << " return tmp;\n"; diff --git a/sc/source/core/opencl/op_spreadsheet.hxx b/sc/source/core/opencl/op_spreadsheet.hxx index 53a0109..0179506 100644 --- a/sc/source/core/opencl/op_spreadsheet.hxx +++ b/sc/source/core/opencl/op_spreadsheet.hxx @@ -20,6 +20,7 @@ public: virtual void GenSlidingWindowFunction(std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments) SAL_OVERRIDE; virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "VLookup"; } + virtual bool takeString() const SAL_OVERRIDE { return true; } }; }} diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx index dfda160..ccd653c 100644 --- a/sc/source/core/opencl/opbase.cxx +++ b/sc/source/core/opencl/opbase.cxx @@ -59,15 +59,15 @@ void VectorRef::GenSlidingWindowDecl(std::stringstream &ss) const } /// When referenced in a sliding window function -std::string VectorRef::GenSlidingWindowDeclRef(bool) const +std::string VectorRef::GenSlidingWindowDeclRef(bool nested) const { std::stringstream ss; formula::SingleVectorRefToken *pSVR = dynamic_cast<formula::SingleVectorRefToken*>(DynamicKernelArgument::GetFormulaToken()); - if (pSVR) + if (pSVR&&nested) ss << "(gid0 < " << pSVR->GetArrayLength() << "?"; ss << mSymName << "[gid0]"; - if (pSVR) + if (pSVR&&nested) ss << ":NAN)"; return ss.str(); } @@ -155,7 +155,7 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ss<< " if("; } ss<< "isNan("; - ss<< vSubArguments[i]->GenSlidingWindowDeclRef(); + ss<< vSubArguments[i]->GenSlidingWindowDeclRef(false); ss<<"))\n"; ss<< " tmp"; ss<< i; @@ -164,7 +164,7 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ss <<" tmp"; ss <<i; ss << "="; - ss << vSubArguments[i]->GenSlidingWindowDeclRef(); + ss << vSubArguments[i]->GenSlidingWindowDeclRef(false); ss<<";\n"; } diff --git a/sc/source/core/opencl/opbase.hxx b/sc/source/core/opencl/opbase.hxx index 6cd04a9..d1d0005 100644 --- a/sc/source/core/opencl/opbase.hxx +++ b/sc/source/core/opencl/opbase.hxx @@ -165,8 +165,14 @@ public: virtual std::string GenStringSlidingWindowDeclRef(bool=false) const { return std::string(""); } + virtual bool IsMixedArgument() const + { return false; } + /// Generate use/references to the argument virtual void GenDeclRef(std::stringstream &ss) const; + virtual void GenNumDeclRef(std::stringstream &ss) const{ss << ",";} + + virtual void GenStringDeclRef(std::stringstream &ss) const{ss << ",";} /// Create buffer and pass the buffer to a given kernel virtual size_t Marshal(cl_kernel, int, int, cl_program) = 0; @@ -205,7 +211,7 @@ public: virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE; /// When referenced in a sliding window function - virtual std::string GenSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE; + virtual std::string GenSlidingWindowDeclRef(bool=true) const SAL_OVERRIDE; /// Create buffer and pass the buffer to a given kernel virtual size_t Marshal(cl_kernel, int, int, cl_program) SAL_OVERRIDE; _______________________________________________ Libreoffice-commits mailing list libreoffice-comm...@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/libreoffice-commits