diff options
author | fengzeng <fengzeng@multicorewareinc.com> | 2013-12-24 11:20:51 +0800 |
---|---|---|
committer | I-Jui (Ray) Sung <ray@multicorewareinc.com> | 2013-12-30 15:05:24 -0600 |
commit | 850f079075b31c9aecd2d6b296c0c5bef6738724 (patch) | |
tree | 26f08072229434fa72bb7873dcb80bdcd9e3932a /sc | |
parent | 4f5f49deb3a6eb549d3950cb4773f0601c1fd52c (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.cxx | 1 | ||||
-rw-r--r-- | sc/source/core/opencl/op_math.cxx | 228 |
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, |