summaryrefslogtreecommitdiff
path: root/sc
diff options
context:
space:
mode:
authorhaochen <haochen@multicorewareinc.com>2014-05-31 13:54:22 +0800
committerMarkus Mohrhard <markus.mohrhard@collabora.co.uk>2014-06-10 15:58:26 +0200
commitc0ea62fcb7f2bd809eea0fd45fb71e31ffba871a (patch)
treeaaa167b44bdec05c741ac152bb6db478a37c37dc /sc
parent0fa58160b90aa9817e5a6fb18c7427895f4b2f4f (diff)
GPU Calc:Support 3rd parameter in FLOOR
Change-Id: Ie3a265f34a5f589d41e802b63df4be6a64989b05
Diffstat (limited to 'sc')
-rw-r--r--sc/source/core/opencl/formulagroupcl.cxx2
-rw-r--r--sc/source/core/opencl/op_math.cxx79
-rw-r--r--sc/source/core/opencl/op_statistical.cxx92
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";