summaryrefslogtreecommitdiff
path: root/sc
diff options
context:
space:
mode:
authorhaochen <haochen@multicorewareinc.com>2014-06-13 11:26:12 +0800
committerMarkus Mohrhard <markus.mohrhard@googlemail.com>2014-06-17 07:43:59 +0200
commit3ad748accd49b128edeb979969d2100af8902624 (patch)
treed3d9cf94184e0a36157da1c11f51995ec31ed7cf /sc
parent8e19f8ebd8febccf7a706f1e0179a4d72f30d6d8 (diff)
GPU Calc:Support nested formulae expansion for simple nested
Change-Id: If1ae42a5481cf76942ff1ac5e0ee31a94159badd
Diffstat (limited to 'sc')
-rw-r--r--sc/source/core/opencl/formulagroupcl.cxx110
-rw-r--r--sc/source/core/opencl/op_logical.cxx54
-rw-r--r--sc/source/core/opencl/opbase.cxx48
-rw-r--r--sc/source/core/opencl/opbase.hxx5
4 files changed, 142 insertions, 75 deletions
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index 07900adc729f..e19d4a2388c1 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 9dce77b73306..96170f2293f3 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 02b9c37bb28d..b0751f53eea5 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 d1d00058a02f..487fc04b09ee 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() {}
};