summaryrefslogtreecommitdiff
path: root/sc
diff options
context:
space:
mode:
authorfengzeng <fengzeng@multicorewareinc.com>2013-12-24 11:20:51 +0800
committerI-Jui (Ray) Sung <ray@multicorewareinc.com>2013-12-30 15:05:24 -0600
commit850f079075b31c9aecd2d6b296c0c5bef6738724 (patch)
tree26f08072229434fa72bb7873dcb80bdcd9e3932a /sc
parent4f5f49deb3a6eb549d3950cb4773f0601c1fd52c (diff)
GPU Calc: Fix bug of SIN,POWER,EXP,SQRT
AMLOEXT-380 Change-Id: I52833dea851f24797fe23e76e1b403716672b091 Signed-off-by: haochen <haochen@multicorewareinc.com> Signed-off-by: Wei Wei <weiwei@multicorewareinc.com> Signed-off-by: I-Jui (Ray) Sung <ray@multicorewareinc.com>
Diffstat (limited to 'sc')
-rw-r--r--sc/source/core/opencl/formulagroupcl.cxx1
-rw-r--r--sc/source/core/opencl/op_math.cxx228
2 files changed, 167 insertions, 62 deletions
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index c5159f0423f9..3264cf4c41d0 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -2228,6 +2228,7 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(
mvSubArguments.push_back(SoPHelper(ts,
ft->Children[i],new OpChiDist));
break;
+ case ocPow:
case ocPower:
mvSubArguments.push_back(SoPHelper(ts,
ft->Children[i], new OpPower));
diff --git a/sc/source/core/opencl/op_math.cxx b/sc/source/core/opencl/op_math.cxx
index 3f742f8f603a..4e1e5f182fa5 100644
--- a/sc/source/core/opencl/op_math.cxx
+++ b/sc/source/core/opencl/op_math.cxx
@@ -826,33 +826,55 @@ void OpCscH::GenSlidingWindowFunction(
ss << "}";
}
void OpExp::GenSlidingWindowFunction(std::stringstream &ss,
- const std::string sSymName, SubArguments &vSubArguments)
+ const std::string sSymName, SubArguments &vSubArguments)
{
- FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken();
- const formula::SingleVectorRefToken*tmpCurDVR= dynamic_cast<const
- formula::SingleVectorRefToken *>(tmpCur);
ss << "\ndouble " << sSymName;
ss << "_"<< BinFuncName() <<"(";
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
- if (i)
- ss << ",";
+ if (i) ss << ",";
vSubArguments[i]->GenSlidingWindowDecl(ss);
}
- ss << ")\n{\n\t";
- ss <<"int gid0=get_global_id(0);\n\t";
- ss << "double arg0 = " << vSubArguments[0]->GenSlidingWindowDeclRef();
- ss << ";\n\t";
+ ss << ")\n";
+ ss << "{\n";
+ ss << " int gid0=get_global_id(0);\n";
+ ss << " double arg0 = 0.0f;\n";
+ FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken();
+ assert(tmpCur);
+ if(ocPush == vSubArguments[0]->GetFormulaToken()->GetOpCode())
+ {
+ if(tmpCur->GetType() == formula::svSingleVectorRef)
+ {
+ const formula::SingleVectorRefToken*tmpCurDVR=
+ dynamic_cast
+ <const formula::SingleVectorRefToken *>(tmpCur);
+ ss << " arg0 = ";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ";\n";
#ifdef ISNAN
- ss<< "if(isNan(arg0)||(gid0>=";
- ss<<tmpCurDVR->GetArrayLength();
- ss<<"))\n\t\t";
- ss<<"arg0 = 0;\n\t";
+ ss << " if(isNan(";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ")||(gid0>=";
+ ss << tmpCurDVR->GetArrayLength();
+ ss << "))\n";
+ ss << " { arg0 = 0.0f; }\n";
#endif
- ss << "double tmp = pow(M_E, arg0);\n\t";
- ss << "return tmp;\n";
+ }
+ else if(tmpCur->GetType() == formula::svDouble)
+ {
+ ss << " arg0=" << tmpCur->GetDouble() << ";\n";
+ }
+ }
+ else
+ {
+ ss << " arg0 = ";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ";\n";
+ }
+ ss << " return pow(M_E, arg0);\n";
ss << "}";
}
+
void OpAverageIfs::GenSlidingWindowFunction(std::stringstream &ss,
const std::string sSymName, SubArguments &vSubArguments)
{
@@ -974,32 +996,54 @@ void OpSinh::GenSlidingWindowFunction(std::stringstream &ss,
void OpSin::GenSlidingWindowFunction(std::stringstream &ss,
const std::string sSymName, SubArguments &vSubArguments)
{
- FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken();
- const formula::SingleVectorRefToken*tmpCurDVR= dynamic_cast<const
- formula::SingleVectorRefToken *>(tmpCur);
ss << "\ndouble " << sSymName;
ss << "_"<< BinFuncName() <<"(";
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
- if (i)
- ss << ",";
+ if (i) ss << ",";
vSubArguments[i]->GenSlidingWindowDecl(ss);
}
ss << ")\n";
ss << "{\n";
ss << " int gid0=get_global_id(0);\n";
- ss << " double arg0 = "<< vSubArguments[0]->GenSlidingWindowDeclRef();
- ss << ";\n";
+ ss << " double arg0 = 0.0f;\n";
+ FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken();
+ assert(tmpCur);
+ if(ocPush == vSubArguments[0]->GetFormulaToken()->GetOpCode())
+ {
+ if(tmpCur->GetType() == formula::svSingleVectorRef)
+ {
+ const formula::SingleVectorRefToken*tmpCurDVR=
+ dynamic_cast
+ <const formula::SingleVectorRefToken *>(tmpCur);
+ ss << " arg0 = ";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ";\n";
#ifdef ISNAN
- ss << " if(isNan(arg0)||(gid0>=";
- ss << tmpCurDVR->GetArrayLength();
- ss << "))\n";
- ss << " arg0 = 0;\n";
+ ss << " if(isNan(";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ")||(gid0>=";
+ ss << tmpCurDVR->GetArrayLength();
+ ss << "))\n";
+ ss << " { arg0 = 0.0f; }\n";
#endif
- ss << " double x = arg0 * M_1_PI;\n";
- ss << " return sinpi(x);\n";
+ }
+ else if(tmpCur->GetType() == formula::svDouble)
+ {
+ ss << " arg0=" << tmpCur->GetDouble() << ";\n";
+ }
+ }
+ else
+ {
+ ss << " arg0 = ";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ";\n";
+ }
+ ss << " arg0 = arg0 * M_1_PI;\n";
+ ss << " return sinpi(arg0);\n";
ss << "}";
}
+
void OpAbs::GenSlidingWindowFunction(std::stringstream &ss,
const std::string sSymName, SubArguments &vSubArguments)
{
@@ -1190,13 +1234,9 @@ void OpTanH::GenSlidingWindowFunction(std::stringstream &ss,
ss << " return tmp;\n";
ss << "}";
}
-void OpPower::GenSlidingWindowFunction(
- std::stringstream &ss, const std::string sSymName,
- SubArguments &vSubArguments)
+void OpPower::GenSlidingWindowFunction(std::stringstream &ss,
+ const std::string sSymName, SubArguments &vSubArguments)
{
- FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken();
- const formula::SingleVectorRefToken*tmpCurDVR= dynamic_cast<const
- formula::SingleVectorRefToken *>(tmpCur);
ss << "\ndouble " << sSymName;
ss << "_"<< BinFuncName() <<"(";
for (unsigned i = 0; i < vSubArguments.size(); i++)
@@ -1205,32 +1245,70 @@ void OpPower::GenSlidingWindowFunction(
ss << ",";
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 << ";\n";
+ ss << ")\n";
+ ss << "{\n";
+ ss << " int gid0=get_global_id(0);\n";
+ ss << " double arg[2];\n";
+ for( unsigned i=0; i < vSubArguments.size(); ++i)
+ {
+ FormulaToken *tmpCur = vSubArguments[i]->GetFormulaToken();
+ assert(tmpCur);
+ if(ocPush == vSubArguments[i]->GetFormulaToken()->GetOpCode())
+ {
+ if(tmpCur->GetType() == formula::svDoubleVectorRef)
+ {
+ const formula::DoubleVectorRefToken* tmpCurDVR =
+ dynamic_cast<
+ const formula::DoubleVectorRefToken *>(tmpCur);
+ ss << " int i = 0;\n";
+ ss << " arg["<<i<<"] = ";
+ ss << vSubArguments[i]->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(arg["<<i;
+ ss << "])||(gid0>=";
+ ss << tmpCurDVR->GetArrayLength();
+ ss << "))\n";
+ ss << " arg["<<i;
+ ss << "] = 0;\n";
#endif
- ss << " double tmp=pow(arg0,arg1);\n";
- ss << " return tmp;\n";
+ }
+ else if(tmpCur->GetType() == formula::svSingleVectorRef)
+ {
+ const formula::SingleVectorRefToken* tmpCurDVR=
+ dynamic_cast<
+ const formula::SingleVectorRefToken *>(tmpCur);
+ ss << " arg["<<i<<"] = ";
+ ss << vSubArguments[i]->GenSlidingWindowDeclRef();
+ ss << ";\n";
+#ifdef ISNAN
+ ss << " if(isNan(arg["<<i;
+ ss << "])||(gid0>=";
+ ss << tmpCurDVR->GetArrayLength();
+ ss << "))\n";
+ ss << " arg["<<i;
+ ss << "] = 0;\n";
+#endif
+ }
+ else if(tmpCur->GetType() == formula::svDouble)
+ {
+ ss << " arg["<<i<<"] = ";
+ ss << tmpCur->GetDouble() << ";\n";
+ }
+ }
+ else
+ {
+ ss << " arg["<<i<<"] = ";
+ ss << vSubArguments[i]->GenSlidingWindowDeclRef();
+ ss << ";\n";
+ }
+ }
+ ss << " return pow(arg[0],arg[1]);\n";
ss << "}";
}
void OpSqrt::GenSlidingWindowFunction(std::stringstream &ss,
const std::string sSymName, SubArguments &vSubArguments)
{
- FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken();
- const formula::SingleVectorRefToken*tmpCurDVR= dynamic_cast<const
- formula::SingleVectorRefToken *>(tmpCur);
ss << "\ndouble " << sSymName;
ss << "_"<< BinFuncName() <<"(";
for (unsigned i = 0; i < vSubArguments.size(); i++)
@@ -1242,16 +1320,42 @@ void OpSqrt::GenSlidingWindowFunction(std::stringstream &ss,
ss << ")\n";
ss << "{\n";
ss << " int gid0=get_global_id(0);\n";
- ss << " double arg0 = "<< vSubArguments[0]->GenSlidingWindowDeclRef();
- ss << ";\n";
+ ss << " double arg0 = 0.0f;\n";
+ FormulaToken *tmpCur = vSubArguments[0]->GetFormulaToken();
+ assert(tmpCur);
+ if(ocPush == vSubArguments[0]->GetFormulaToken()->GetOpCode())
+ {
+ if(tmpCur->GetType() == formula::svSingleVectorRef)
+ {
+ const formula::SingleVectorRefToken*tmpCurDVR=
+ dynamic_cast
+ <const formula::SingleVectorRefToken *>(tmpCur);
+ ss << " arg0 = ";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ";\n";
#ifdef ISNAN
- ss << " if(isNan(arg0)||(gid0>=";
- ss << tmpCurDVR->GetArrayLength();
- ss << "))\n";
- ss << " arg0 = 0;\n";
+ ss << " if(isNan(";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ")||(gid0>=";
+ ss << tmpCurDVR->GetArrayLength();
+ ss << "))\n";
+ ss << " { arg0 = 0; }\n";
#endif
- ss << " double tmp=sqrt(arg0);\n";
- ss << " return tmp;\n";
+ }
+ else if(tmpCur->GetType() == formula::svDouble)
+ {
+ printf("const\n");
+ ss << " arg0=";
+ ss << tmpCur->GetDouble() << ";\n";
+ }
+ }
+ else
+ {
+ ss << " arg0 = ";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef();
+ ss << ";\n";
+ }
+ ss << " return sqrt(arg0);\n";
ss << "}";
}
void OpArcCot::GenSlidingWindowFunction(std::stringstream &ss,