GPU Calc: optimize countifs

Change-Id: Id378cf89b0742b23c8282148223cb4a7d4317aac
Signed-off-by: I-Jui (Ray) Sung <ray@multicorewareinc.com>

sc/source/core/opencl/op_math.cxx

@@ -394,48 +394,38 @@ void OpCountIfs::GenSlidingWindowFunction(std::stringstream &ss,
     }
     ss << ")\n    {\n";
     ss <<"    int gid0=get_global_id(0);\n";
-    ss << "    double tmp =0;\n";
-    ss << "    int i ;\n";
+    ss << "    int tmp =0;\n";
+    ss << "    int loop;\n";
     GenTmpVariables(ss,vSubArguments);
-    ss << "    for (i = ";
-     if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) {
-        ss << "gid0; i < "<< nCurWindowSize <<"; i++)\n";
-     } else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) {
-        ss << "0; i < gid0+"<< nCurWindowSize <<"; i++)\n";
-     } else {
-        ss << "0; i < "<< nCurWindowSize <<"; i++)\n";
-     }
-     ss << "    {\n";
-     if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
-     {
-        ss<< "    int doubleIndex =i+gid0;\n";
-     }else
-     {
-        ss<< "    int doubleIndex =i;\n";
-     }
+
      ss<< "    int singleIndex =gid0;\n";
      int m=0;
+
+     std::stringstream tmpss;
+
      for(unsigned j=0;j<vSubArguments.size();j+=2,m++)
      {
-        CheckSubArgumentIsNan(ss,vSubArguments,j);
+        CheckSubArgumentIsNan(tmpss,vSubArguments,j);
         CheckSubArgumentIsNan(ss,vSubArguments,j+1);
-        ss <<"    if(isequal(";
-        ss <<"tmp";
-        ss <<j;
-        ss <<" , ";
-        ss << "tmp";
-        ss << j+1;
-        ss << ")){\n";
+        tmpss <<"    if(isequal(";
+        tmpss <<"tmp";
+        tmpss <<j;
+        tmpss <<" , ";
+        tmpss << "tmp";
+        tmpss << j+1;
+        tmpss << ")){\n";
      }
-    ss << "    tmp =tmp +1;\n";
-    for(unsigned j=0;j<=vSubArguments.size();j+=2,m--)
+    tmpss << "    tmp ++;\n";
+    for(unsigned j=0;j<vSubArguments.size();j+=2,m--)
      {
          for(int n = 0;n<m+1;n++)
         {
-            ss << "    ";
+            tmpss << "    ";
         }
-        ss<< "}\n";
+        tmpss<< "}\n";
      }
+     UnrollDoubleVector(ss,tmpss,pCurDVR,nCurWindowSize);
+
     ss << "return tmp;\n";
     ss << "}";
 }

sc/source/core/opencl/opbase.cxx

@@ -9,8 +9,6 @@
 
 #include "opbase.hxx"
 
-#include "formula/vectortoken.hxx"
-
 using namespace formula;
 
 namespace sc { namespace opencl {
@@ -224,6 +222,69 @@ void CheckVariables::CheckAllSubArgumentIsNan(
     }
 }
 
+void CheckVariables::UnrollDoubleVector( std::stringstream & ss,
+std::stringstream & unrollstr,const formula::DoubleVectorRefToken* pCurDVR,
+int nCurWindowSize)
+{
+    int unrollSize = 16;
+    if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) {
+        ss << "    loop = ("<<nCurWindowSize<<" - gid0)/";
+        ss << unrollSize<<";\n";
+    } else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) {
+        ss << "    loop = ("<<nCurWindowSize<<" + gid0)/";
+        ss << unrollSize<<";\n";
+
+    } else {
+        ss << "    loop = "<<nCurWindowSize<<"/"<< unrollSize<<";\n";
+    }
+
+    ss << "    for ( int j = 0;j< loop; j++)\n";
+    ss << "    {\n";
+    ss << "        int i = ";
+    if (!pCurDVR->IsStartFixed()&& pCurDVR->IsEndFixed()) {
+       ss << "gid0 + j * "<< unrollSize <<";\n";
+    }else {
+       ss << "j * "<< unrollSize <<";\n";
+    }
+
+    if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
+    {
+       ss << "        int doubleIndex = i+gid0;\n";
+    }else
+    {
+       ss << "        int doubleIndex = i;\n";
+    }
+
+    for(int j =0;j < unrollSize;j++)
+    {
+        ss << unrollstr.str();
+        ss << "i++;\n";
+        ss << "doubleIndex++;\n";
+    }
+     ss << "    }\n";
+     ss << "    for (int i = ";
+     if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) {
+        ss << "gid0 + loop *"<<unrollSize<<"; i < ";
+        ss << nCurWindowSize <<"; i++)\n";
+     } else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) {
+        ss << "0 + loop *"<<unrollSize<<"; i < gid0+";
+        ss << nCurWindowSize <<"; i++)\n";
+     } else {
+        ss << "0 + loop *"<<unrollSize<<"; i < ";
+        ss << nCurWindowSize <<"; i++)\n";
+     }
+     ss << "    {\n";
+     if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
+     {
+        ss << "        int doubleIndex = i+gid0;\n";
+     }else
+     {
+        ss << "        int doubleIndex = i;\n";
+     }
+     ss << unrollstr.str();
+     ss << "    }\n";
+}
+
 }}
 
 /* vim:set shiftwidth=4 softtabstop=4 expandtab: */

sc/source/core/opencl/opbase.hxx

@@ -15,7 +15,7 @@
 #include "clcc/clew.h"
 
 #include "formula/token.hxx"
-
+#include "formula/vectortoken.hxx"
 #include <boost/shared_ptr.hpp>
 #include <boost/noncopyable.hpp>
 #include <set>
@@ -266,6 +266,9 @@ public:
     // only check isNan
     void CheckSubArgumentIsNan2(std::stringstream &ss,
             SubArguments &vSubArguments, int argumentNum, std::string p);
+    void UnrollDoubleVector(std::stringstream &ss,
+    std::stringstream &unrollstr, const formula::DoubleVectorRefToken* pCurDVR,
+    int nCurWindowSize);
 };
 
 }}