diff options
author | haochen <haochen@multicorewareinc.com> | 2014-05-31 13:54:22 +0800 |
---|---|---|
committer | Markus Mohrhard <markus.mohrhard@collabora.co.uk> | 2014-06-10 15:58:26 +0200 |
commit | c0ea62fcb7f2bd809eea0fd45fb71e31ffba871a (patch) | |
tree | aaa167b44bdec05c741ac152bb6db478a37c37dc /sc | |
parent | 0fa58160b90aa9817e5a6fb18c7427895f4b2f4f (diff) |
GPU Calc:Support 3rd parameter in FLOOR
Change-Id: Ie3a265f34a5f589d41e802b63df4be6a64989b05
Diffstat (limited to 'sc')
-rw-r--r-- | sc/source/core/opencl/formulagroupcl.cxx | 2 | ||||
-rw-r--r-- | sc/source/core/opencl/op_math.cxx | 79 | ||||
-rw-r--r-- | sc/source/core/opencl/op_statistical.cxx | 92 |
3 files changed, 79 insertions, 94 deletions
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index 91a54f4b81c1..9b45c9124d60 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 4a960b36a79d..36c94f62a099 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 88541e263913..20dbb49362ba 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"; |