summaryrefslogtreecommitdiff
path: root/sc/source/core/opencl/op_math.cxx
diff options
context:
space:
mode:
authoryangzhang <yangzhang@multicorewareinc.com>2013-11-12 09:27:25 +0800
committerI-Jui (Ray) Sung <ray@multicorewareinc.com>2013-11-15 11:54:25 -0600
commit61c3a61baa2054f2962b922614f0098a203c33a3 (patch)
treec901b83222eafbf02e199c74e8cf7a6b11880be2 /sc/source/core/opencl/op_math.cxx
parent5fd9b52dd48d932ce652c9ef1084e033ea755493 (diff)
GPU Calc: implemented CEIL
AMLOEXT-181 FIX Change-Id: I3ee0eb38e071ec0b0224fc73a47cdd0408b345d4 Signed-off-by: haochen <haochen@multicorewareinc.com> Signed-off-by: I-Jui (Ray) Sung <ray@multicorewareinc.com>
Diffstat (limited to 'sc/source/core/opencl/op_math.cxx')
-rw-r--r--sc/source/core/opencl/op_math.cxx72
1 files changed, 72 insertions, 0 deletions
diff --git a/sc/source/core/opencl/op_math.cxx b/sc/source/core/opencl/op_math.cxx
index 78e42497289e..b3198624d2a4 100644
--- a/sc/source/core/opencl/op_math.cxx
+++ b/sc/source/core/opencl/op_math.cxx
@@ -1720,6 +1720,78 @@ void OpSqrtPi::GenSlidingWindowFunction(std::stringstream &ss,
ss << " 3.1415926535897932384626433832795f);\n";
ss << "}";
}
+void OpCeil::GenSlidingWindowFunction(std::stringstream &ss,
+ const std::string sSymName, SubArguments &vSubArguments)
+{
+ ss << "\ndouble " << sSymName;
+ ss << "_"<< BinFuncName() <<"(";
+ for (unsigned i = 0; i < vSubArguments.size(); i++)
+ {
+ if (i) ss << ",";
+ vSubArguments[i]->GenSlidingWindowDecl(ss);
+ }
+ ss << ") {\n";
+ ss << " int gid0 = get_global_id(0);\n";
+ ss << " double num = " << GetBottom() << ";\n";
+ ss << " double significance = " << GetBottom() << ";\n";
+ ss << " double bAbs = 0;\n";
+#ifdef ISNAN
+ FormulaToken *iNum = vSubArguments[0]->GetFormulaToken();
+ const formula::SingleVectorRefToken* tmpCurDVRNum=
+ dynamic_cast<const formula::SingleVectorRefToken*>(iNum);
+ FormulaToken *iSignificance = vSubArguments[1]->GetFormulaToken();
+ const formula::SingleVectorRefToken* tmpCurDVRSignificance=
+ dynamic_cast<const formula::SingleVectorRefToken*>(iSignificance);
+ ss << " int buffer_num_len = "<<tmpCurDVRNum->GetArrayLength() << ";\n";
+ ss << " int buffer_significance_len = ";
+ ss << tmpCurDVRSignificance->GetArrayLength() << ";\n";
+ ss << " if((gid0)>=buffer_num_len || isNan(";
+ ss << vSubArguments[0]->GenSlidingWindowDeclRef() << "))\n";
+ ss << " num = " << GetBottom() << ";\n";
+ ss << " else\n ";
+#endif
+ ss << " num = " << vSubArguments[0]->GenSlidingWindowDeclRef() << ";\n";
+#ifdef ISNAN
+ ss << " if((gid0)>=buffer_significance_len || isNan(";
+ ss << vSubArguments[1]->GenSlidingWindowDeclRef() << "))\n";
+ ss << " significance = " << GetBottom() << ";\n";
+ ss << " else\n ";
+#endif
+ ss << " significance = ";
+ ss << vSubArguments[1]->GenSlidingWindowDeclRef() << ";\n";
+ if (vSubArguments.size() > 2)
+ {
+#ifdef ISNAN
+ FormulaToken *bAbs = vSubArguments[2]->GetFormulaToken();
+ if(bAbs->GetType() == formula::svSingleVectorRef)
+ {
+ const formula::SingleVectorRefToken* tmpCurSVRIsAbs=
+ dynamic_cast<const formula::SingleVectorRefToken*>(bAbs);
+ ss<< " if((gid0)>=" << tmpCurSVRIsAbs->GetArrayLength() << " ||";
+ }
+ if(bAbs->GetType() == formula::svDoubleVectorRef)
+ {
+ const formula::DoubleVectorRefToken* tmpCurDVRIsAbs=
+ dynamic_cast<const formula::DoubleVectorRefToken*>(bAbs);
+ ss<< " if((gid0)>=" << tmpCurDVRIsAbs->GetArrayLength() << " ||";
+ }
+ if(bAbs->GetType() == formula::svDouble)
+ {
+ ss<< " if(";
+ }
+ ss << "isNan(";
+ ss << vSubArguments[2]->GenSlidingWindowDeclRef() << "))\n";
+ ss << " bAbs = 0;\n";
+ ss << " else\n ";
+#endif
+ ss << " bAbs = "<<vSubArguments[2]->GenSlidingWindowDeclRef()<<";\n";
+ }
+ ss << " return ";
+ ss << "( !(int)bAbs && num < 0.0 ? floor( num / significance ) : ";
+ ss << "ceil( num / significance ) )";
+ ss << "*significance;\n";
+ ss << "}";
+}
}}
/* vim:set shiftwidth=4 softtabstop=4 expandtab: */