Skip to content

C
core
Kaydet (Commit) 3ad748ac authored tarafından haochen's avatar haochen Kaydeden (comit) Markus Mohrhard
Dosyalara gözat
Seçenekler

GPU Calc:Support nested formulae expansion for simple nested 

Change-Id: If1ae42a5481cf76942ff1ac5e0ee31a94159badd
üst 8e19f8eb
Show whitespace changes
Inline Side-by-side
Showing with 115 additions and 48 deletions
sc/source/core/opencl/formulagroupcl.cxx
Dosyayı görüntüle @ 3ad748ac
...@@ -1256,12 +1256,12 @@ public: ...@@ -1256,12 +1256,12 @@ public:
ss << ", "; ss << ", ";
vSubArguments[i]->GenSlidingWindowDecl(ss); vSubArguments[i]->GenSlidingWindowDecl(ss);
} }
ss << ") {\n\t"; ss << ") {\n";
ss << "double tmp = " << GetBottom() <<";\n\t"; ss << "double tmp = " << GetBottom() <<";\n";
ss << "int gid0 = get_global_id(0);\n\t"; ss << "int gid0 = get_global_id(0);\n";
if (isAverage()) if (isAverage())
ss << "int nCount = 0;\n\t"; ss << "int nCount = 0;\n";
ss << "double tmpBottom;\n\t"; ss << "double tmpBottom;\n";
unsigned i = vSubArguments.size(); unsigned i = vSubArguments.size();
while (i--) while (i--)
{ {
...@@ -1292,60 +1292,52 @@ public: ...@@ -1292,60 +1292,52 @@ public:
if (pCur->GetType() == formula::svSingleVectorRef) if (pCur->GetType() == formula::svSingleVectorRef)
{ {
#ifdef ISNAN
const formula::SingleVectorRefToken* pSVR = const formula::SingleVectorRefToken* pSVR =
static_cast< const formula::SingleVectorRefToken* >(pCur); static_cast< const formula::SingleVectorRefToken* >(pCur);
ss << "if (gid0 < " << pSVR->GetArrayLength() << "){\n\t\t"; ss << "if (gid0 < " << pSVR->GetArrayLength() << "){\n";
#else
#endif
} }
else if (pCur->GetType() == formula::svDouble) else if (pCur->GetType() == formula::svDouble)
{ {
#ifdef ISNAN ss << "{\n";
ss << "{\n\t\t";
#endif
}
else
{
} }
} }
#ifdef ISNAN
if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode()) if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
{ {
ss << "tmpBottom = " << GetBottom() << ";\n\t\t"; ss << "tmpBottom = " << GetBottom() << ";\n";
ss << "if (isNan("; ss << "if (isNan(";
ss << vSubArguments[i]->GenSlidingWindowDeclRef(); ss << vSubArguments[i]->GenSlidingWindowDeclRef();
ss << "))\n\t\t\t"; ss << "))\n";
ss << "tmp = "; if ( ZeroReturnZero() )
ss << Gen2("tmpBottom", "tmp") << ";\n\t\t"; ss << " return 0;\n";
ss << "else{\n\t\t\t"; else
ss << "tmp = "; {
ss << " tmp = ";
ss << Gen2("tmpBottom", "tmp") << ";\n";
}
ss << "else{\n";
ss << " tmp = ";
ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp"); ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp");
ss << ";\n\t\t\t"; ss << ";\n";
ss << "}\n\t"; ss << " }\n";
ss << "}\n\t"; ss << "}\n";
if ( vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svSingleVectorRef&& ZeroReturnZero() )
{
ss << "else{\n";
ss << " return 0;\n";
ss << " }\n";
}
} }
else else
{ {
ss << "tmp = "; ss << "tmp = ";
ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "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"; ss << "return tmp";
#ifdef ISNAN
if (isAverage()) if (isAverage())
ss << "*pow((double)nCount,-1.0)"; ss << "*pow((double)nCount,-1.0)";
#else
if (isAverage())
ss << "/(double)"<<nItems;
#endif
ss << ";\n}"; ss << ";\n}";
} }
virtual bool isAverage() const { return false; } virtual bool isAverage() const { return false; }
...@@ -1535,6 +1527,11 @@ public: ...@@ -1535,6 +1527,11 @@ public:
static_cast< const formula::SingleVectorRefToken*> static_cast< const formula::SingleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken()); (vSubArguments[i]->GetFormulaToken());
temp3<<pSVR->GetArrayLength(); temp3<<pSVR->GetArrayLength();
temp3 << ")||isNan("<<vSubArguments[i]
->GenSlidingWindowDeclRef();
temp3 << ")?0:";
temp3 << vSubArguments[i]->GenSlidingWindowDeclRef();
temp3 << ")";
} }
else if(vSubArguments[i]->GetFormulaToken()->GetType() == else if(vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svDoubleVectorRef){ formula::svDoubleVectorRef){
...@@ -1542,13 +1539,14 @@ public: ...@@ -1542,13 +1539,14 @@ public:
static_cast< const formula::DoubleVectorRefToken*> static_cast< const formula::DoubleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken()); (vSubArguments[i]->GetFormulaToken());
temp3<<pSVR->GetArrayLength(); temp3<<pSVR->GetArrayLength();
}
temp3 << ")||isNan("<<vSubArguments[i] temp3 << ")||isNan("<<vSubArguments[i]
->GenSlidingWindowDeclRef(true); ->GenSlidingWindowDeclRef(true);
temp3 << ")?0:"; temp3 << ")?0:";
temp3 << vSubArguments[i]->GenSlidingWindowDeclRef(true); temp3 << vSubArguments[i]->GenSlidingWindowDeclRef(true);
temp3 << ")"; temp3 << ")";
} }
}
else else
temp3 << vSubArguments[i]->GenSlidingWindowDeclRef(true); temp3 << vSubArguments[i]->GenSlidingWindowDeclRef(true);
} }
...@@ -1604,6 +1602,11 @@ public: ...@@ -1604,6 +1602,11 @@ public:
static_cast< const formula::SingleVectorRefToken*> static_cast< const formula::SingleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken()); (vSubArguments[i]->GetFormulaToken());
temp4<<pSVR->GetArrayLength(); temp4<<pSVR->GetArrayLength();
temp4 << ")||isNan("<<vSubArguments[i]
->GenSlidingWindowDeclRef();
temp4 << ")?0:";
temp4 << vSubArguments[i]->GenSlidingWindowDeclRef();
temp4 << ")";
} }
else if(vSubArguments[i]->GetFormulaToken()->GetType() == else if(vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svDoubleVectorRef) formula::svDoubleVectorRef)
...@@ -1612,13 +1615,14 @@ public: ...@@ -1612,13 +1615,14 @@ public:
static_cast< const formula::DoubleVectorRefToken*> static_cast< const formula::DoubleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken()); (vSubArguments[i]->GetFormulaToken());
temp4<<pSVR->GetArrayLength(); temp4<<pSVR->GetArrayLength();
}
temp4 << ")||isNan("<<vSubArguments[i] temp4 << ")||isNan("<<vSubArguments[i]
->GenSlidingWindowDeclRef(true); ->GenSlidingWindowDeclRef(true);
temp4 << ")?0:"; temp4 << ")?0:";
temp4 << vSubArguments[i]->GenSlidingWindowDeclRef(true); temp4 << vSubArguments[i]->GenSlidingWindowDeclRef(true);
temp4 << ")"; temp4 << ")";
} }
}
else else
{ {
temp4 << vSubArguments[i] temp4 << vSubArguments[i]
...@@ -1752,6 +1756,7 @@ public: ...@@ -1752,6 +1756,7 @@ public:
return lhs + "*" + rhs; return lhs + "*" + rhs;
} }
virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "fmul"; } virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "fmul"; }
virtual bool ZeroReturnZero() {return true;};
}; };
/// Technically not a reduction, but fits the framework. /// Technically not a reduction, but fits the framework.
...@@ -2021,8 +2026,19 @@ public: ...@@ -2021,8 +2026,19 @@ public:
} else { } else {
if (mvSubArguments.size() != 2) if (mvSubArguments.size() != 2)
throw Unhandled(); throw Unhandled();
ss << "(" << mpCodeGen->Gen2(mvSubArguments[0]->GenSlidingWindowDeclRef(true), bool bArgument1_NeedNested =
mvSubArguments[1]->GenSlidingWindowDeclRef(true)) << ")"; (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(); return ss.str();
} }
......
sc/source/core/opencl/op_logical.cxx
Dosyayı görüntüle @ 3ad748ac
...@@ -319,7 +319,6 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss, ...@@ -319,7 +319,6 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss,
{ {
ss << "\ndouble " << sSymName; ss << "\ndouble " << sSymName;
ss << "_"<< BinFuncName() <<"("; ss << "_"<< BinFuncName() <<"(";
if(vSubArguments.size()!=3) throw Unhandled("unknown operand for ocPush");
for (unsigned i = 0; i < vSubArguments.size(); i++) for (unsigned i = 0; i < vSubArguments.size(); i++)
{ {
if (i) if (i)
...@@ -335,6 +334,8 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss, ...@@ -335,6 +334,8 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss,
throw UnhandledToken(tmpCur0, "unknown operand for ocPush"); throw UnhandledToken(tmpCur0, "unknown operand for ocPush");
} }
else else
{
if(vSubArguments.size()==3)
{ {
ss << " if(isNan("; ss << " if(isNan(";
ss << vSubArguments[0]->GenSlidingWindowDeclRef(); ss << vSubArguments[0]->GenSlidingWindowDeclRef();
...@@ -349,6 +350,31 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss, ...@@ -349,6 +350,31 @@ void OpIf::GenSlidingWindowFunction(std::stringstream &ss,
ss << vSubArguments[1]->GenSlidingWindowDeclRef(); ss << vSubArguments[1]->GenSlidingWindowDeclRef();
ss <<";\n"; 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"; ss << "}\n";
} }
......
sc/source/core/opencl/opbase.cxx
Dosyayı görüntüle @ 3ad748ac
...@@ -130,7 +130,6 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ...@@ -130,7 +130,6 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss,
SubArguments &vSubArguments, int argumentNum) SubArguments &vSubArguments, int argumentNum)
{ {
int i = argumentNum; int i = argumentNum;
#ifdef ISNAN
if(vSubArguments[i]->GetFormulaToken()->GetType() == if(vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svSingleVectorRef) formula::svSingleVectorRef)
{ {
...@@ -139,6 +138,17 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ...@@ -139,6 +138,17 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss,
ss<< " if(singleIndex>="; ss<< " if(singleIndex>=";
ss<< pTmpDVR1->GetArrayLength(); ss<< pTmpDVR1->GetArrayLength();
ss<<" ||"; 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() == if(vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svDoubleVectorRef) formula::svDoubleVectorRef)
...@@ -148,24 +158,36 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss, ...@@ -148,24 +158,36 @@ void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss,
ss<< " if(doubleIndex>="; ss<< " if(doubleIndex>=";
ss<< pTmpDVR2->GetArrayLength(); ss<< pTmpDVR2->GetArrayLength();
ss<<" ||"; 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 || if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble ||
vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush) vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush)
{ {
ss<< " if("; ss<< " if(";
}
ss<< "isNan("; ss<< "isNan(";
ss<< vSubArguments[i]->GenSlidingWindowDeclRef(true); ss<< vSubArguments[i]->GenSlidingWindowDeclRef();
ss<<"))\n"; ss<<"))\n";
ss<< " tmp"; ss<< " tmp";
ss<< i; ss<< i;
ss <<"=0;\n else \n"; ss <<"=0;\n else \n";
#endif
ss <<" tmp"; ss <<" tmp";
ss <<i; ss <<i;
ss << "="; ss << "=";
ss << vSubArguments[i]->GenSlidingWindowDeclRef(true); ss << vSubArguments[i]->GenSlidingWindowDeclRef();
ss<<";\n"; ss<<";\n";
}
} }
void CheckVariables::CheckSubArgumentIsNan2( std::stringstream & ss, void CheckVariables::CheckSubArgumentIsNan2( std::stringstream & ss,
......
sc/source/core/opencl/opbase.hxx
Dosyayı görüntüle @ 3ad748ac
...@@ -211,7 +211,7 @@ public: ...@@ -211,7 +211,7 @@ public:
virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE; virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE;
/// When referenced in a sliding window function /// 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 /// Create buffer and pass the buffer to a given kernel
virtual size_t Marshal(cl_kernel, int, int, cl_program) SAL_OVERRIDE; virtual size_t Marshal(cl_kernel, int, int, cl_program) SAL_OVERRIDE;
...@@ -250,6 +250,9 @@ public: ...@@ -250,6 +250,9 @@ public:
std::set<std::string>& ) {} std::set<std::string>& ) {}
virtual bool takeString() const = 0; virtual bool takeString() const = 0;
virtual bool takeNumeric() 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() {} virtual ~OpBase() {}
}; };
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment