summaryrefslogtreecommitdiff
path: root/sc
diff options
context:
space:
mode:
authorhaochen <haochen@multicorewareinc.com>2014-05-13 09:35:08 +0800
committerMarkus Mohrhard <markus.mohrhard@googlemail.com>2014-05-19 09:13:22 +0200
commitfbbf5c089d9e102f819df752dce0e1b3223ad269 (patch)
treeec041e9ae7390bd827ce5195d2f8226b5dd08ad9 /sc
parentdd0f844728a53a337233a3c687b792e725803da5 (diff)
GPU Calc: Fixed style&error problem in formulagroupcl and op_addin
Change-Id: Id7c6e341eb88f6c8ffdaa7f83b787bb19649afec
Diffstat (limited to 'sc')
-rw-r--r--sc/source/core/opencl/formulagroupcl.cxx31
-rw-r--r--sc/source/core/opencl/op_addin.cxx65
-rw-r--r--sc/source/core/opencl/op_addin.hxx7
3 files changed, 88 insertions, 15 deletions
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index 7066c4fc3ba8..4f40e2567d00 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -31,7 +31,7 @@
/// CONFIGURATIONS
// Comment out this to turn off FMIN and FMAX intrinsics
#define USE_FMIN_FMAX 1
-#define REDUCE_THRESHOLD 4 // set to 4 for correctness testing. priority 1
+#define REDUCE_THRESHOLD 201 // set to 4 for correctness testing. priority 1
#define UNROLLING_FACTOR 16 // set to 4 for correctness testing (if no reduce)
#include "formulagroupcl_public.hxx"
#ifdef WIN32
@@ -1218,30 +1218,25 @@ public:
ss << "int nCount = 0;\n\t";
ss << "double tmpBottom;\n\t";
unsigned i = vSubArguments.size();
- size_t nItems = 0;
while (i--)
{
if (NumericRange *NR =
dynamic_cast<NumericRange *> (vSubArguments[i].get()))
{
- bool needBody;
- nItems += NR->GenReductionLoopHeader(ss, needBody);
- if (needBody == false) continue;
+ bool needBody;NR->GenReductionLoopHeader(ss, needBody);if (needBody == false) continue;
}
else if (ParallelNumericRange *PNR =
dynamic_cast<ParallelNumericRange *> (vSubArguments[i].get()))
{
//did not handle yet
- bool needBody;
- nItems += PNR->GenReductionLoopHeader(ss, needBody);
- if (needBody == false) continue;
+ bool needBody;PNR->GenReductionLoopHeader(ss, needBody);if (needBody == false) continue;
}
else if (StringRange *SR =
dynamic_cast<StringRange *> (vSubArguments[i].get()))
{
//did not handle yet
bool needBody;
- nItems += SR->GenReductionLoopHeader(ss, needBody);
+ SR->GenReductionLoopHeader(ss, needBody);
if (needBody == false) continue;
}
else
@@ -1257,7 +1252,6 @@ public:
static_cast< const formula::SingleVectorRefToken* >(pCur);
ss << "if (gid0 < " << pSVR->GetArrayLength() << "){\n\t\t";
#else
- nItems += 1;
#endif
}
else if (pCur->GetType() == formula::svDouble)
@@ -1265,11 +1259,9 @@ public:
#ifdef ISNAN
ss << "{\n\t\t";
#endif
- nItems += 1;
}
else
{
- nItems += 1;
}
}
#ifdef ISNAN
@@ -1304,7 +1296,7 @@ public:
ss << "return tmp";
#ifdef ISNAN
if (isAverage())
- ss << "/(double)nCount";
+ ss << "*pow((double)nCount,-1.0)";
#else
if (isAverage())
ss << "/(double)"<<nItems;
@@ -2827,10 +2819,10 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(
mvSubArguments.push_back(SoPHelper(ts,
ft->Children[i], new OpSumIf));
break;
- /*case ocNegSub:
+ case ocNegSub:
mvSubArguments.push_back(SoPHelper(ts,
ft->Children[i], new OpNegSub));
- break;*/
+ break;
case ocAveDev:
mvSubArguments.push_back(SoPHelper(ts,
ft->Children[i], new OpAveDev));
@@ -3071,6 +3063,12 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(
mvSubArguments.push_back(SoPHelper(ts, ft->Children[i],
new OpBesselj));
}
+ else if ( !(pChild->GetExternal().compareTo(OUString(
+ "com.sun.star.sheet.addin.Analysis.getGestep"))))
+ {
+ mvSubArguments.push_back(SoPHelper(ts, ft->Children[i],
+ new OpGestep));
+ }
else
throw UnhandledToken(pChild, "unhandled opcode");
break;
@@ -3175,6 +3173,9 @@ public:
global_work_size, NULL, 0, NULL, NULL);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
+ err = clFinish(kEnv.mpkCmdQueue);
+ if (CL_SUCCESS != err)
+ throw OpenCLError(err, __FILE__, __LINE__);
}
virtual ~DynamicKernel();
cl_mem GetResultBuffer(void) const { return mpResClmem; }
diff --git a/sc/source/core/opencl/op_addin.cxx b/sc/source/core/opencl/op_addin.cxx
index 06c4fbb8b542..1cc2fe1f027a 100644
--- a/sc/source/core/opencl/op_addin.cxx
+++ b/sc/source/core/opencl/op_addin.cxx
@@ -196,6 +196,71 @@ void OpBesselj::GenSlidingWindowFunction(std::stringstream &ss,
ss << " return DBL_MAX;\n";
ss << "}";
}
+void OpGestep::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 << "{\n";
+ ss << " double tmp=0,tmp0 =0,tmp1 = 0;\n";
+ ss << " int gid0=get_global_id(0);\n";
+ size_t i = vSubArguments.size();
+ ss <<"\n";
+ for (i = 0; i < vSubArguments.size(); i++)
+ {
+ FormulaToken *pCur = vSubArguments[i]->GetFormulaToken();
+ assert(pCur);
+ if (pCur->GetType() == formula::svSingleVectorRef)
+ {
+#ifdef ISNAN
+ const formula::SingleVectorRefToken* pSVR =
+ dynamic_cast< const formula::SingleVectorRefToken* >(pCur);
+ ss << " if (gid0 < " << pSVR->GetArrayLength() << ")\n";
+ ss << " {\n";
+#endif
+ }
+ else if (pCur->GetType() == formula::svDouble)
+ {
+#ifdef ISNAN
+ ss << " {\n";
+#endif
+ }
+ else
+ {
+#ifdef ISNAN
+#endif
+ }
+#ifdef ISNAN
+ if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
+ {
+ ss << " if (isNan(";
+ ss << vSubArguments[i]->GenSlidingWindowDeclRef();
+ ss << "))\n";
+ ss << " tmp"<<i<<" = 0;\n";
+ ss << " else\n";
+ ss << " tmp"<<i<<" = ";
+ ss << vSubArguments[i]->GenSlidingWindowDeclRef();
+ ss << ";\n }\n";
+ }
+ else
+ {
+ ss << "tmp"<<i<<" ="<<vSubArguments[i]->GenSlidingWindowDeclRef();
+ ss <<";\n";
+ }
+#endif
+ }
+ ss << " tmp =tmp0 >= tmp1 ? 1 : 0;\n";
+ ss << " return tmp;\n";
+ ss << "}\n";
+}
}}
diff --git a/sc/source/core/opencl/op_addin.hxx b/sc/source/core/opencl/op_addin.hxx
index 56bc868a4b18..ecdec028392b 100644
--- a/sc/source/core/opencl/op_addin.hxx
+++ b/sc/source/core/opencl/op_addin.hxx
@@ -21,6 +21,13 @@ public:
const std::string &sSymName, SubArguments &vSubArguments) SAL_OVERRIDE;
virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "Besselj"; }
};
+class OpGestep: public Normal
+{
+public:
+ virtual void GenSlidingWindowFunction(std::stringstream &ss,
+ const std::string &sSymName, SubArguments &vSubArguments) SAL_OVERRIDE;;
+ virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "Gestep"; }
+};
}}