[Libreoffice-commits] core.git: 7 commits - config_host/config_features.h.in configure.ac sc/inc sc/Library_sc.mk sc/source

Michael Meeks michael.meeks at suse.com
Thu Jun 27 11:33:54 PDT 2013


 config_host/config_features.h.in         |    7 
 configure.ac                             |    1 
 sc/Library_sc.mk                         |    1 
 sc/inc/document.hxx                      |    1 
 sc/inc/formulagroup.hxx                  |   30 
 sc/source/core/data/documen2.cxx         |    2 
 sc/source/core/data/formulacell.cxx      |    4 
 sc/source/core/opencl/formulagroupcl.cxx |  287 +++++
 sc/source/core/opencl/oclkernels.hxx     |  144 +-
 sc/source/core/opencl/openclwrapper.cxx  | 1517 +++++++++++++++++++------------
 sc/source/core/opencl/openclwrapper.hxx  |  160 +--
 sc/source/core/tool/formulagroup.cxx     |   62 -
 sc/source/core/tool/interpr1.cxx         |    4 
 sc/source/ui/app/scmod.cxx               |    9 
 sc/source/ui/optdlg/calcoptionsdlg.cxx   |    4 
 15 files changed, 1498 insertions(+), 735 deletions(-)

New commits:
commit 66e72c38f77bedc28b990b2a724829b50c70bf8d
Author: Michael Meeks <michael.meeks at suse.com>
Date:   Thu Jun 27 16:49:41 2013 +0100

    abstract out the FormulaGroupInterpreter more cleanly.

diff --git a/sc/Library_sc.mk b/sc/Library_sc.mk
index 94c9a52..e96bb14 100644
--- a/sc/Library_sc.mk
+++ b/sc/Library_sc.mk
@@ -58,6 +58,7 @@ ifeq ($(ENABLE_OPENCL),TRUE)
 $(eval $(call gb_Library_use_externals,sc,opencl))
 
 $(eval $(call gb_Library_add_exception_objects,sc,\
+	sc/source/core/opencl/formulagroupcl \
 	sc/source/core/opencl/openclwrapper \
 ))
 endif
diff --git a/sc/inc/document.hxx b/sc/inc/document.hxx
index ff0b869..377c80e 100644
--- a/sc/inc/document.hxx
+++ b/sc/inc/document.hxx
@@ -415,6 +415,7 @@ private:
     ::std::set<ScFormulaCell*> maSubTotalCells;
 
     bool                mbUseEmbedFonts;
+
 public:
     bool              IsUsingEmbededFonts() { return mbUseEmbedFonts; }
     void              SetIsUsingEmbededFonts( bool bUse ) { mbUseEmbedFonts = bUse; }
diff --git a/sc/inc/formulagroup.hxx b/sc/inc/formulagroup.hxx
index 55f6e09..e5839e4 100644
--- a/sc/inc/formulagroup.hxx
+++ b/sc/inc/formulagroup.hxx
@@ -30,19 +30,31 @@ struct FormulaGroupContext : boost::noncopyable
 };
 
 /**
- * All the vectorized formula calculation code should be collectd here.
+ * All the vectorized formula calculation code should be collected here.
+ *
+ * Abstract base class for formula group interpreters, and a factory.
  */
-class FormulaGroupInterpreter
+class SC_DLLPUBLIC FormulaGroupInterpreter
+{
+    static FormulaGroupInterpreter *msInstance;
+ protected:
+    FormulaGroupInterpreter() {}
+    virtual ~FormulaGroupInterpreter() {}
+ public:
+    static FormulaGroupInterpreter *getStatic();
+
+    virtual bool interpret(ScDocument& rDoc, const ScAddress& rTopPos, const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode) = 0;
+};
+
+/// Inherit from this for alternate formula group calculation approaches.
+class SC_DLLPUBLIC FormulaGroupInterpreterSoftware : public FormulaGroupInterpreter
 {
-    ScDocument& mrDoc;
-    ScAddress maTopPos;
-    ScFormulaCellGroupRef mxGroup;
-    ScTokenArray& mrCode;
 public:
-    FormulaGroupInterpreter(
-        ScDocument& rDoc, const ScAddress& rTopPos, const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode);
+    FormulaGroupInterpreterSoftware() :
+        FormulaGroupInterpreter() {}
+    virtual ~FormulaGroupInterpreterSoftware() {}
 
-    bool interpret();
+    virtual bool interpret(ScDocument& rDoc, const ScAddress& rTopPos, const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode);
 };
 
 }
diff --git a/sc/source/core/data/documen2.cxx b/sc/source/core/data/documen2.cxx
index b1b3160..6244eac 100644
--- a/sc/source/core/data/documen2.cxx
+++ b/sc/source/core/data/documen2.cxx
@@ -86,6 +86,8 @@
 #include "macromgr.hxx"
 #include "formulacell.hxx"
 #include "clipcontext.hxx"
+#include "interpre.hxx"
+#include "formulagroup.hxx"
 
 using namespace com::sun::star;
 
diff --git a/sc/source/core/data/formulacell.cxx b/sc/source/core/data/formulacell.cxx
index 68d242d..f57da6f 100644
--- a/sc/source/core/data/formulacell.cxx
+++ b/sc/source/core/data/formulacell.cxx
@@ -3094,9 +3094,7 @@ bool ScFormulaCell::InterpretFormulaGroup()
     GroupTokenConverter aConverter(aCode, *pDocument, *this);
     if (!aConverter.convert(*pCode))
         return false;
-
-    sc::FormulaGroupInterpreter aInterpreter(*pDocument, aPos, xGroup, aCode);
-    return aInterpreter.interpret();
+    return sc::FormulaGroupInterpreter::getStatic()->interpret(*pDocument, aPos, xGroup, aCode);
 }
 
 bool ScFormulaCell::InterpretInvariantFormulaGroup()
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
new file mode 100644
index 0000000..54fd8e8
--- /dev/null
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -0,0 +1,287 @@
+/* -*- Mode: C++; tab-width: 4; indent-tabs-mode: nil; c-basic-offset: 4 -*- */
+/*
+ * This file is part of the LibreOffice project.
+ *
+ * This Source Code Form is subject to the terms of the Mozilla Public
+ * License, v. 2.0. If a copy of the MPL was not distributed with this
+ * file, You can obtain one at http://mozilla.org/MPL/2.0/.
+ */
+
+#include <config_features.h>
+#include "formulagroup.hxx"
+#include "document.hxx"
+#include "formulacell.hxx"
+#include "tokenarray.hxx"
+#include "compiler.hxx"
+#include "interpre.hxx"
+#include "formula/vectortoken.hxx"
+
+#include "openclwrapper.hxx"
+
+namespace sc {
+
+// A single public entry point for a factory function:
+namespace opencl {
+    extern sc::FormulaGroupInterpreter *createFormulaGroupInterpreter();
+}
+
+/////time test dbg
+double getTimeDiff(const TimeValue& t1, const TimeValue& t2)
+{
+    double tv1 = t1.Seconds;
+    double tv2 = t2.Seconds;
+    tv1 += t1.Nanosec / 1000000000.0;
+    tv2 += t2.Nanosec / 1000000000.0;
+
+    return tv1 - tv2;
+}//dbg-t
+TimeValue aTimeBefore, aTimeAfter;
+///////////////////////////////////////
+
+class FormulaGroupInterpreterOpenCL : public FormulaGroupInterpreterSoftware
+{
+public:
+    FormulaGroupInterpreterOpenCL() :
+        FormulaGroupInterpreterSoftware()
+    {
+        OclCalc::InitEnv();
+    }
+    virtual ~FormulaGroupInterpreterOpenCL()
+    {
+        OclCalc::ReleaseOpenclRunEnv();
+    }
+    virtual bool interpret(ScDocument& rDoc, const ScAddress& rTopPos,
+                           const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode);
+};
+
+bool FormulaGroupInterpreterOpenCL::interpret(ScDocument& rDoc, const ScAddress& rTopPos,
+                                              const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode)
+{
+    size_t rowSize = xGroup->mnLength; //, srcSize = 0;
+    fprintf(stderr,"rowSize at begin is ...%ld.\n",(long)rowSize);
+    int *rangeStart =NULL; // The first position for calculation,for example,the A1 in (=MAX(A1:A100))
+    int *rangeEnd = NULL; // The last position for calculation,for example, the A100 in (=MAX(A1:A100))
+    // The row quantity can be gotten from p2->GetArrayLength()
+    int count1 =0,count2 =0,count3=0;
+    int oclOp=0;
+    double *srcData = NULL; // Point to the input data from CPU
+    double *rResult=NULL; // Point to the output data from GPU
+    double *leftData=NULL; // Left input for binary operator(+,-,*,/),for example,(=leftData+rightData)
+    double *rightData=NULL; // Right input for binary operator(+,-,*,/),for example,(=leftData/rightData)
+                            // The rightData can't be zero for "/"
+
+    leftData  = (double *)malloc(sizeof(double) * rowSize);
+    rightData = (double *)malloc(sizeof(double) * rowSize);
+    rResult   = (double *)malloc(sizeof(double) * rowSize*2);// For 2 columns(B,C)
+    srcData = (double *)calloc(rowSize,sizeof(double));
+
+    rangeStart =(int *)malloc(sizeof(int) * rowSize);
+    rangeEnd   =(int *)malloc(sizeof(int) * rowSize);
+
+    memset(rResult,0,rowSize);
+    if(NULL==leftData||NULL==rightData||
+           NULL==rResult||NULL==rangeStart||NULL==rangeEnd)
+    {
+        printf("malloc err\n");
+        return false;
+    }
+    // printf("rowSize is %d.\n",rowsize);
+
+    // Until we implement group calculation for real, decompose the group into
+    // individual formula token arrays for individual calculation.
+    ScAddress aTmpPos = rTopPos;
+    for (sal_Int32 i = 0; i < xGroup->mnLength; ++i)
+    {
+        aTmpPos.SetRow(xGroup->mnStart + i);
+        ScTokenArray aCode2;
+        for (const formula::FormulaToken* p = rCode.First(); p; p = rCode.Next())
+        {
+            switch (p->GetType())
+            {
+                case formula::svSingleVectorRef:
+                {
+                    const formula::SingleVectorRefToken* p2 = static_cast<const formula::SingleVectorRefToken*>(p);
+                    const double* pArray = p2->GetArray();
+                    aCode2.AddDouble(static_cast<size_t>(i) < p2->GetArrayLength() ? pArray[i] : 0.0);
+                }
+                break;
+                case formula::svDoubleVectorRef:
+                {
+                    const formula::DoubleVectorRefToken* p2 = static_cast<const formula::DoubleVectorRefToken*>(p);
+                    const std::vector<const double*>& rArrays = p2->GetArrays();
+                    size_t nColSize = rArrays.size();
+                    size_t nRowStart = p2->IsStartFixed() ? 0 : i;
+                    size_t nRowEnd = p2->GetRefRowSize() - 1;
+                    if (!p2->IsEndFixed())
+                        nRowEnd += i;
+                    size_t nRowSize = nRowEnd - nRowStart + 1;
+                    ScMatrixRef pMat(new ScMatrix(nColSize, nRowSize, 0.0));
+
+                    //srcSize = rowSize+nRowSize-rowSize%nRowSize;//align as nRowSize
+                    //srcData = (double *)calloc(srcSize,sizeof(double));
+                    rangeStart[i] = nRowStart;//record the start position
+                    rangeEnd[i] = nRowEnd;//record the end position
+
+                    for (size_t nCol = 0; nCol < nColSize; ++nCol)
+                    {
+                        const double* pArray = rArrays[nCol];
+
+                        //printf("pArray is %p.\n",pArray);
+                        if( NULL==pArray )
+                        {
+                            fprintf(stderr,"Error: pArray is NULL!\n");
+                            return false;
+                        }
+                        //fprintf(stderr,"(rowSize+nRowSize-1) is %d.\n",rowSize+nRowSize-1);
+                        for( size_t u=0; u<rowSize; u++ )
+                        {
+                            srcData[u] = pArray[u];// note:rowSize<=srcSize
+                            //fprintf(stderr,"srcData[%d] is %f.\n",u,srcData[u]);
+                        }
+
+                        for (size_t nRow = 0; nRow < nRowSize; ++nRow)
+                        {
+                            if (nRowStart + nRow < p2->GetArrayLength())
+                            {
+                                double fVal = pArray[nRowStart+nRow];
+                                pMat->PutDouble(fVal, nCol, nRow);
+                            }
+                        }
+                    }
+
+                    ScMatrixToken aTok(pMat);
+                    aCode2.AddToken(aTok);
+                }
+                break;
+                default:
+                    aCode2.AddToken(*p);
+            }
+        }
+
+        ScFormulaCell* pDest = rDoc.GetFormulaCell(aTmpPos);
+        if (!pDest)
+            return false;
+
+        const formula::FormulaToken *pCur = aCode2.First();
+        aCode2.Reset();
+        while( ( pCur = aCode2.Next() ) != NULL )
+        {
+            OpCode eOp = pCur->GetOpCode();
+            if(eOp==0)
+            {
+                  if(count3%2==0)
+                    leftData[count1++] = pCur->GetDouble();
+                   else
+                    rightData[count2++] = pCur->GetDouble();
+                count3++;
+               }
+               else if( eOp!=ocOpen && eOp!=ocClose )
+                oclOp = eOp;
+
+//            if(count1>0){//dbg
+//                fprintf(stderr,"leftData is %f.\n",leftData[count1-1]);
+//                count1--;
+//            }
+//            if(count2>0){//dbg
+//                fprintf(stderr,"rightData is %f.\n",rightData[count2-1]);
+//                count2--;
+//            }
+        }
+
+        if(!getenv("SC_GPU"))
+        {
+            fprintf(stderr,"ccCPU flow...\n\n");
+            ScCompiler aComp(&rDoc, aTmpPos, aCode2);
+            aComp.SetGrammar(rDoc.GetGrammar());
+            aComp.CompileTokenArray(); // Create RPN token array.
+            ScInterpreter aInterpreter(pDest, &rDoc, aTmpPos, aCode2);
+            aInterpreter.Interpret();
+            pDest->SetResultToken(aInterpreter.GetResultToken().get());
+            pDest->ResetDirty();
+            pDest->SetChanged(true);
+        }
+    } // for loop end (xGroup->mnLength)
+
+    // For GPU calculation
+    if(getenv("SC_GPU"))
+    {
+            fprintf(stderr,"ggGPU flow...\n\n");
+            printf(" oclOp is... %d\n",oclOp);
+            osl_getSystemTime(&aTimeBefore); //timer
+            static OclCalc ocl_calc;
+            switch(oclOp)
+            {
+                case ocAdd:
+                       ocl_calc.OclHostSignedAdd(leftData,rightData,rResult,count1);
+                    break;
+                case ocSub:
+                    ocl_calc.OclHostSignedSub(leftData,rightData,rResult,count1);
+                    break;
+                case ocMul:
+                    ocl_calc.OclHostSignedMul(leftData,rightData,rResult,count1);
+                    break;
+                case ocDiv:
+                    ocl_calc.OclHostSignedDiv(leftData,rightData,rResult,count1);
+                    break;
+                case ocMax:
+                    ocl_calc.OclHostFormulaMax(srcData,rangeStart,rangeEnd,rResult,rowSize);
+                    break;
+                case ocMin:
+                    ocl_calc.OclHostFormulaMin(srcData,rangeStart,rangeEnd,rResult,rowSize);
+                    break;
+                case ocAverage:
+                    ocl_calc.OclHostFormulaAverage(srcData,rangeStart,rangeEnd,rResult,rowSize);
+                    break;
+                default:
+                    fprintf(stderr,"No OpenCL function for this calculation.\n");
+                    break;
+            }
+            /////////////////////////////////////////////////////
+            osl_getSystemTime(&aTimeAfter);
+            double diff = getTimeDiff(aTimeAfter, aTimeBefore);
+            //if (diff >= 1.0)
+            {
+                fprintf(stderr,"OpenCL,diff...%f.\n",diff);
+            }
+/////////////////////////////////////////////////////
+
+//rResult[i];
+//            for(sal_Int32 i = 0; i < rowSize; ++i){//dbg output results
+//                fprintf(stderr,"After GPU,rRsults[%d] is ...%f\n",i,rResult[i]);
+//            }
+
+            // Insert the double data, in rResult[i] back into the document
+            rDoc.SetFormulaResults(rTopPos, rResult, xGroup->mnLength);
+        }
+
+        if(leftData)
+            free(leftData);
+        if(rightData)
+            free(rightData);
+        if(rangeStart)
+            free(rangeStart);
+        if(rangeEnd)
+            free(rangeEnd);
+        if(rResult)
+            free(rResult);
+        if(srcData)
+            free(srcData);
+
+        if(getenv("SC_GPUSAMPLE")){
+            //fprintf(stderr,"FormulaGroupInterpreter::interpret(),iniflag...%d\n",ocl_calc.GetOpenclState());
+            //ocl_calc.OclTest();//opencl test sample for debug
+        }
+
+    return true;
+}
+
+namespace opencl {
+    sc::FormulaGroupInterpreter *createFormulaGroupInterpreter()
+    {
+        return new sc::FormulaGroupInterpreterOpenCL();
+    }
+} // namespace opencl
+
+} // namespace sc
+
+/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/inc/openclwrapper.hxx b/sc/source/core/opencl/openclwrapper.hxx
similarity index 100%
rename from sc/source/core/inc/openclwrapper.hxx
rename to sc/source/core/opencl/openclwrapper.hxx
diff --git a/sc/source/core/tool/formulagroup.cxx b/sc/source/core/tool/formulagroup.cxx
index 9d9b504..f9d251b 100644
--- a/sc/source/core/tool/formulagroup.cxx
+++ b/sc/source/core/tool/formulagroup.cxx
@@ -16,70 +16,20 @@
 #include "interpre.hxx"
 #include "formula/vectortoken.hxx"
 
-#if HAVE_FEATURE_OPENCL
-#  include "openclwrapper.hxx"
-#endif
-
 namespace sc {
 
-FormulaGroupInterpreter::FormulaGroupInterpreter(
-    ScDocument& rDoc, const ScAddress& rTopPos, const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode) :
-    mrDoc(rDoc), maTopPos(rTopPos), mxGroup(xGroup), mrCode(rCode) {}
-
-/////time test dbg
-double getTimeDiff(const TimeValue& t1, const TimeValue& t2)
-{
-    double tv1 = t1.Seconds;
-    double tv2 = t2.Seconds;
-    tv1 += t1.Nanosec / 1000000000.0;
-    tv2 += t2.Nanosec / 1000000000.0;
-
-    return tv1 - tv2;
-}//dbg-t
-TimeValue aTimeBefore, aTimeAfter;
-///////////////////////////////////////
-
-bool FormulaGroupInterpreter::interpret()
+bool FormulaGroupInterpreterSoftware::interpret(ScDocument& rDoc, const ScAddress& rTopPos,
+                                                const ScFormulaCellGroupRef& xGroup,
+                                                ScTokenArray& rCode)
 {
-#if HAVE_FEATURE_OPENCL
-    size_t rowSize = mxGroup->mnLength; //, srcSize = 0;
-    fprintf(stderr,"rowSize at begin is ...%ld.\n",(long)rowSize);
-    int *rangeStart =NULL; // The first position for calculation,for example,the A1 in (=MAX(A1:A100))
-    int *rangeEnd = NULL; // The last position for calculation,for example, the A100 in (=MAX(A1:A100))
-    // The row quantity can be gotten from p2->GetArrayLength()
-    int count1 =0,count2 =0,count3=0;
-    int oclOp=0;
-    double *srcData = NULL; // Point to the input data from CPU
-    double *rResult=NULL; // Point to the output data from GPU
-    double *leftData=NULL; // Left input for binary operator(+,-,*,/),for example,(=leftData+rightData)
-    double *rightData=NULL; // Right input for binary operator(+,-,*,/),for example,(=leftData/rightData)
-                            // The rightData can't be zero for "/"
-
-    leftData  = (double *)malloc(sizeof(double) * rowSize);
-    rightData = (double *)malloc(sizeof(double) * rowSize);
-    rResult   = (double *)malloc(sizeof(double) * rowSize*2);// For 2 columns(B,C)
-    srcData = (double *)calloc(rowSize,sizeof(double));
-
-    rangeStart =(int *)malloc(sizeof(int) * rowSize);
-    rangeEnd   =(int *)malloc(sizeof(int) * rowSize);
-
-    memset(rResult,0,rowSize);
-    if(NULL==leftData||NULL==rightData||
-           NULL==rResult||NULL==rangeStart||NULL==rangeEnd)
-    {
-        printf("malloc err\n");
-        return false;
-    }
-    // printf("rowSize is %d.\n",rowsize);
-#endif
     // Until we implement group calculation for real, decompose the group into
     // individual formula token arrays for individual calculation.
-    ScAddress aTmpPos = maTopPos;
-    for (sal_Int32 i = 0; i < mxGroup->mnLength; ++i)
+    ScAddress aTmpPos = rTopPos;
+    for (sal_Int32 i = 0; i < xGroup->mnLength; ++i)
     {
-        aTmpPos.SetRow(mxGroup->mnStart + i);
+        aTmpPos.SetRow(xGroup->mnStart + i);
         ScTokenArray aCode2;
-        for (const formula::FormulaToken* p = mrCode.First(); p; p = mrCode.Next())
+        for (const formula::FormulaToken* p = rCode.First(); p; p = rCode.Next())
         {
             switch (p->GetType())
             {
@@ -101,29 +51,9 @@ bool FormulaGroupInterpreter::interpret()
                         nRowEnd += i;
                     size_t nRowSize = nRowEnd - nRowStart + 1;
                     ScMatrixRef pMat(new ScMatrix(nColSize, nRowSize, 0.0));
-#if HAVE_FEATURE_OPENCL
-                    //srcSize = rowSize+nRowSize-rowSize%nRowSize;//align as nRowSize
-                    //srcData = (double *)calloc(srcSize,sizeof(double));
-                    rangeStart[i] = nRowStart;//record the start position
-                    rangeEnd[i] = nRowEnd;//record the end position
-#endif
                     for (size_t nCol = 0; nCol < nColSize; ++nCol)
                     {
                         const double* pArray = rArrays[nCol];
-#if HAVE_FEATURE_OPENCL
-                        //printf("pArray is %p.\n",pArray);
-                        if( NULL==pArray )
-                        {
-                            fprintf(stderr,"Error: pArray is NULL!\n");
-                            return false;
-                        }
-                        //fprintf(stderr,"(rowSize+nRowSize-1) is %d.\n",rowSize+nRowSize-1);
-                        for( size_t u=0; u<rowSize; u++ )
-                        {
-                            srcData[u] = pArray[u];// note:rowSize<=srcSize
-                            //fprintf(stderr,"srcData[%d] is %f.\n",u,srcData[u]);
-                        }
-#endif
                         for (size_t nRow = 0; nRow < nRowSize; ++nRow)
                         {
                             if (nRowStart + nRow < p2->GetArrayLength())
@@ -143,123 +73,53 @@ bool FormulaGroupInterpreter::interpret()
             }
         }
 
-        ScFormulaCell* pDest = mrDoc.GetFormulaCell(aTmpPos);
+        ScFormulaCell* pDest = rDoc.GetFormulaCell(aTmpPos);
         if (!pDest)
             return false;
 
-#if HAVE_FEATURE_OPENCL
-        const formula::FormulaToken *pCur = aCode2.First();
-        aCode2.Reset();
-        while( ( pCur = aCode2.Next() ) != NULL )
-        {
-            OpCode eOp = pCur->GetOpCode();
-            if(eOp==0)
-            {
-                  if(count3%2==0)
-                    leftData[count1++] = pCur->GetDouble();
-                   else
-                    rightData[count2++] = pCur->GetDouble();
-                count3++;
-               }
-               else if( eOp!=ocOpen && eOp!=ocClose )
-                oclOp = eOp;
+        ScCompiler aComp(&rDoc, aTmpPos, aCode2);
+        aComp.SetGrammar(rDoc.GetGrammar());
+        aComp.CompileTokenArray(); // Create RPN token array.
+        ScInterpreter aInterpreter(pDest, &rDoc, aTmpPos, aCode2);
+        aInterpreter.Interpret();
+        pDest->SetResultToken(aInterpreter.GetResultToken().get());
+        pDest->ResetDirty();
+        pDest->SetChanged(true);
+    } // for loop end (xGroup->mnLength)
 
-//            if(count1>0){//dbg
-//                fprintf(stderr,"leftData is %f.\n",leftData[count1-1]);
-//                count1--;
-//            }
-//            if(count2>0){//dbg
-//                fprintf(stderr,"rightData is %f.\n",rightData[count2-1]);
-//                count2--;
-//            }
-        }
-#endif
-        if(!getenv("SC_GPU"))
-        {
-            fprintf(stderr,"ccCPU flow...\n\n");
-            ScCompiler aComp(&mrDoc, aTmpPos, aCode2);
-            aComp.SetGrammar(mrDoc.GetGrammar());
-            aComp.CompileTokenArray(); // Create RPN token array.
-            ScInterpreter aInterpreter(pDest, &mrDoc, aTmpPos, aCode2);
-            aInterpreter.Interpret();
-            pDest->SetResultToken(aInterpreter.GetResultToken().get());
-            pDest->ResetDirty();
-            pDest->SetChanged(true);
-        }
-    } // for loop end (mxGroup->mnLength)
-    // For GPU calculation
-#if HAVE_FEATURE_OPENCL //dbg: Using "export SC_GPU=1" to open if{} in terminal
-    if(getenv("SC_GPU"))
-    {
-            fprintf(stderr,"ggGPU flow...\n\n");
-            printf(" oclOp is... %d\n",oclOp);
-            osl_getSystemTime(&aTimeBefore); //timer
-            static OclCalc ocl_calc;
-            switch(oclOp)
-            {
-                case ocAdd:
-                       ocl_calc.OclHostSignedAdd(leftData,rightData,rResult,count1);
-                    break;
-                case ocSub:
-                    ocl_calc.OclHostSignedSub(leftData,rightData,rResult,count1);
-                    break;
-                case ocMul:
-                    ocl_calc.OclHostSignedMul(leftData,rightData,rResult,count1);
-                    break;
-                case ocDiv:
-                    ocl_calc.OclHostSignedDiv(leftData,rightData,rResult,count1);
-                    break;
-                case ocMax:
-                    ocl_calc.OclHostFormulaMax(srcData,rangeStart,rangeEnd,rResult,rowSize);
-                    break;
-                case ocMin:
-                    ocl_calc.OclHostFormulaMin(srcData,rangeStart,rangeEnd,rResult,rowSize);
-                    break;
-                case ocAverage:
-                    ocl_calc.OclHostFormulaAverage(srcData,rangeStart,rangeEnd,rResult,rowSize);
-                    break;
-                default:
-                    fprintf(stderr,"No OpenCL function for this calculation.\n");
-                    break;
-            }
-            /////////////////////////////////////////////////////
-            osl_getSystemTime(&aTimeAfter);
-            double diff = getTimeDiff(aTimeAfter, aTimeBefore);
-            //if (diff >= 1.0)
-            {
-                fprintf(stderr,"OpenCL,diff...%f.\n",diff);
-            }
-/////////////////////////////////////////////////////
+    return true;
+}
 
-//rResult[i];
-//            for(sal_Int32 i = 0; i < rowSize; ++i){//dbg output results
-//                fprintf(stderr,"After GPU,rRsults[%d] is ...%f\n",i,rResult[i]);
-//            }
+// TODO: load module, hook symbol out, check it works, UI on failure etc.
+namespace opencl {
+    extern sc::FormulaGroupInterpreter *createFormulaGroupInterpreter();
+}
 
-            // Insert the double data, in rResult[i] back into the document
-            mrDoc.SetFormulaResults(maTopPos, rResult, mxGroup->mnLength);
-        }
+FormulaGroupInterpreter *FormulaGroupInterpreter::msInstance = NULL;
 
-        if(leftData)
-            free(leftData);
-        if(rightData)
-            free(rightData);
-        if(rangeStart)
-            free(rangeStart);
-        if(rangeEnd)
-            free(rangeEnd);
-        if(rResult)
-            free(rResult);
-        if(srcData)
-            free(srcData);
+/// load and/or configure the correct formula group interpreter
+FormulaGroupInterpreter *FormulaGroupInterpreter::getStatic()
+{
+    static bool bOpenCLEnabled = false;
 
-        if(getenv("SC_GPUSAMPLE")){
-            //fprintf(stderr,"FormulaGroupInterpreter::interpret(),iniflag...%d\n",ocl_calc.GetOpenclState());
-            //ocl_calc.OclTest();//opencl test sample for debug
-        }
+    if ( msInstance &&
+         bOpenCLEnabled != ScInterpreter::GetGlobalConfig().mbOpenCLEnabled )
+    {
+        delete msInstance;
+        msInstance = NULL;
+    }
+
+    if ( !msInstance )
+    {
+#if HAVE_FEATURE_OPENCL
+        if ( ScInterpreter::GetGlobalConfig().mbOpenCLEnabled )
+            msInstance = sc::opencl::createFormulaGroupInterpreter();
 #endif
+        if ( !msInstance ) // software fallback
+            msInstance = new sc::FormulaGroupInterpreterSoftware();
+    }
 
-    return true;
+    return msInstance;
 }
 
 }
diff --git a/sc/source/ui/app/scmod.cxx b/sc/source/ui/app/scmod.cxx
index 4edfcf0..d9b5d02 100644
--- a/sc/source/ui/app/scmod.cxx
+++ b/sc/source/ui/app/scmod.cxx
@@ -102,10 +102,7 @@
 
 #include "scabstdlg.hxx"
 #include "formula/errorcodes.hxx"
-
-#if HAVE_FEATURE_OPENCL
-#  include "openclwrapper.hxx"
-#endif
+#include "formulagroup.hxx"
 
 #define SC_IDLE_MIN     150
 #define SC_IDLE_MAX     3000
@@ -154,12 +151,8 @@ ScModule::ScModule( SfxObjectFactory* pFact ) :
     mbIsInSharedDocLoading( false ),
     mbIsInSharedDocSaving( false )
 {
-#if HAVE_FEATURE_OPENCL
-    OclCalc::InitEnv();
-#endif
     //  im ctor ist der ResManager (DLL-Daten) noch nicht initialisiert!
-
-    SetName(OUString("StarCalc"));       // fuer Basic
+    SetName(OUString("StarCalc"));       // for Basic
 
     ResetDragObject();
     SetClipObject( NULL, NULL );
@@ -186,13 +179,13 @@ ScModule::ScModule( SfxObjectFactory* pFact ) :
     ScGlobal::InitTextHeight( pMessagePool );
 
     StartListening( *SFX_APP() );       // for SFX_HINT_DEINITIALIZING
+
+    // initialize formula grouping
+    sc::FormulaGroupInterpreter::getStatic();
 }
 
 ScModule::~ScModule()
 {
-#if HAVE_FEATURE_OPENCL
-    OclCalc::ReleaseOpenclRunEnv();
-#endif
     OSL_ENSURE( !pSelTransfer, "Selection Transfer object not deleted" );
 
     //  InputHandler braucht nicht mehr geloescht zu werden (gibt keinen an der App mehr)
commit d7f2e338cf00052f7edbecf664b3d0088198a374
Author: Michael Meeks <michael.meeks at suse.com>
Date:   Wed Jun 26 22:15:19 2013 +0100

    return is a statement not a function.

diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 4790510..20d60e2 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -152,20 +152,20 @@ int OpenclDevice::ConvertToString(const char *filename, char **source) {
         rewind(file);
         *source = (char*) malloc(sizeof(char) * file_size + 1);
         if (*source == (char*) NULL) {
-            return (0);
+            return 0;
         }
         result = fread(*source, 1, file_size, file);
         if (result != (size_t) file_size) {
             free(*source);
-            return (0);
+            return 0;
         }
         (*source)[file_size] = '\0';
         fclose(file);
 
-        return (1);
+        return 1;
     }
     printf("open kernel file failed.\n");
-    return (0);
+    return 0;
 }
 
 int OpenclDevice::BinaryGenerated(const char * clFileName, FILE ** fhandle) {
@@ -388,12 +388,12 @@ int OpenclDevice::CachedOfKernerPrg(const GPUEnv *gpuEnvCached,
     for (i = 0; i < gpuEnvCached->fileCount; i++) {
         if (strcasecmp(gpuEnvCached->kernelSrcFile[i], clFileName) == 0) {
             if (gpuEnvCached->programs[i] != NULL) {
-                return (1);
+                return 1;
             }
         }
     }
 
-    return (0);
+    return 0;
 }
 
 int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) {
@@ -409,7 +409,7 @@ int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) {
     const char* filename = "kernel.cl";
 	fprintf(stderr, "CompileKernelFile ... \n");
     if (CachedOfKernerPrg(gpuInfo, filename) == 1) {
-        return (1);
+        return 1;
     }
 
     idx = gpuInfo->fileCount;
@@ -471,7 +471,7 @@ int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) {
     }
 
     if (gpuInfo->programs[idx] == (cl_program) NULL) {
-        return (0);
+        return 0;
     }
 
     //char options[512];
@@ -498,11 +498,11 @@ int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) {
         }
         if (status != CL_SUCCESS) {
             printf("opencl create build log fail\n");
-            return (0);
+            return 0;
         }
         buildLog = (char*) malloc(length);
         if (buildLog == (char*) NULL) {
-            return (0);
+            return 0;
         }
         if (!gpuInfo->isUserCreated) {
             status = clGetProgramBuildInfo(gpuInfo->programs[idx],
@@ -521,7 +521,7 @@ int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) {
         }
 
         free(buildLog);
-        return (0);
+        return 0;
     }
 
     strcpy(gpuEnv.kernelSrcFile[idx], filename);
@@ -531,7 +531,7 @@ int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) {
 
     gpuInfo->fileCount += 1;
 
-    return (1);
+    return 1;
 
 
 }
@@ -546,10 +546,10 @@ int OpenclDevice::GetKernelEnvAndFunc(const char *kernelName,
             env->program = gpuEnv.programs[0];
             env->kernel = gpuEnv.kernels[i];
             *function = gpuEnv.kernelFunctions[i];
-            return (1);
+            return 1;
         }
     }
-    return (0);
+    return 0;
 }
 
 int OpenclDevice::RunKernel(const char *kernelName, void **userdata) {
@@ -565,11 +565,11 @@ int OpenclDevice::RunKernel(const char *kernelName, void **userdata) {
     if (status == 1) {
         if (&env == (KernelEnv *) NULL
                 || &function == (cl_kernel_function *) NULL) {
-            return (0);
+            return 0;
         }
         return (function(userdata, &env));
     }
-    return (0);
+    return 0;
 }
 int OpenclDevice::InitOpenclRunEnv(int argc, const char *buildOptionKernelfiles)
 {
@@ -587,19 +587,19 @@ int OpenclDevice::InitOpenclRunEnv(int argc, const char *buildOptionKernelfiles)
         status = InitOpenclRunEnv(&gpuEnv);
         if (status) {
             printf("init_opencl_env failed.\n");
-            return (1);
+            return 1;
         }
         printf("init_opencl_env successed.\n");
         //initialize program, kernelName, kernelCount
         status = CompileKernelFile( &gpuEnv, buildOptionKernelfiles);
         if (status == 0 || gpuEnv.kernelCount == 0) {
             printf("CompileKernelFile failed.\n");
-            return (1);
+            return 1;
         }
         printf("CompileKernelFile successed.\n");
         isInited = 1;
     }
-    return (0);
+    return 0;
 }
 
 int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
@@ -617,7 +617,7 @@ int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
     if (!gpuInfo->isUserCreated) {
         status = clGetPlatformIDs(0, NULL, &numPlatforms);
         if (status != CL_SUCCESS) {
-            return (1);
+            return 1;
         }
         gpuInfo->platform = NULL;
 
@@ -625,12 +625,12 @@ int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
             platforms = (cl_platform_id*) malloc(
                     numPlatforms * sizeof(cl_platform_id));
             if (platforms == (cl_platform_id*) NULL) {
-                return (1);
+                return 1;
             }
             status = clGetPlatformIDs(numPlatforms, platforms, NULL);
 
             if (status != CL_SUCCESS) {
-                return (1);
+                return 1;
             }
 
             for (i = 0; i < numPlatforms; i++) {
@@ -638,7 +638,7 @@ int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
                         sizeof(platformName), platformName, NULL);
 
                 if (status != CL_SUCCESS) {
-                    return (1);
+                    return 1;
                 }
                 gpuInfo->platform = platforms[i];
 
@@ -654,7 +654,7 @@ int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
 												&numDevices);
 
                     if (status != CL_SUCCESS) {
-                        return (1);
+                        return 1;
                     }
 
                     if (numDevices) {
@@ -665,7 +665,7 @@ int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
             free(platforms);
         }
         if (NULL == gpuInfo->platform) {
-            return (1);
+            return 1;
         }
 
         // Use available platform.
@@ -692,25 +692,25 @@ int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
         }
         if ((gpuInfo->context == (cl_context) NULL)
                 || (status != CL_SUCCESS)) {
-            return (1);
+            return 1;
         }
         // Detect OpenCL devices.
         // First, get the size of device list data
         status = clGetContextInfo(gpuInfo->context, CL_CONTEXT_DEVICES, 0,
                 NULL, &length);
         if ((status != CL_SUCCESS) || (length == 0)) {
-            return (1);
+            return 1;
         }
         // Now allocate memory for device list based on the size we got earlier
         gpuInfo->devices = (cl_device_id*) malloc(length);
         if (gpuInfo->devices == (cl_device_id*) NULL) {
-            return (1);
+            return 1;
         }
         // Now, get the device list data
         status = clGetContextInfo(gpuInfo->context, CL_CONTEXT_DEVICES, length,
                 gpuInfo->devices, NULL);
         if (status != CL_SUCCESS) {
-            return (1);
+            return 1;
         }
 
         // Create OpenCL command queue.
@@ -718,7 +718,7 @@ int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
                 gpuInfo->devices[0], 0, &status);
 
         if (status != CL_SUCCESS) {
-            return (1);
+            return 1;
         }
     }
 
@@ -737,10 +737,10 @@ int OpenclDevice::RegisterKernelWrapper(const char *kernelName,cl_kernel_functio
 		if (strcasecmp(kernelName, gpuEnv.kernelNames[i]) == 0)
 		{
 			gpuEnv.kernelFunctions[i] = function;
-			return (1);
+			return 1;
 		}
 	}
-    return (0);
+    return 0;
 }
 
 
commit 5cedd85bdccd6275ba146aaa68ecd3d63d60c60c
Author: Michael Meeks <michael.meeks at suse.com>
Date:   Wed Jun 26 22:08:09 2013 +0100

    Use new method for returning values.

diff --git a/sc/source/core/tool/formulagroup.cxx b/sc/source/core/tool/formulagroup.cxx
index 2262cc5..9d9b504 100644
--- a/sc/source/core/tool/formulagroup.cxx
+++ b/sc/source/core/tool/formulagroup.cxx
@@ -42,8 +42,8 @@ TimeValue aTimeBefore, aTimeAfter;
 bool FormulaGroupInterpreter::interpret()
 {
 #if HAVE_FEATURE_OPENCL
-    size_t rowSize = mxGroup->mnLength, srcSize = 0;
-    fprintf(stderr,"rowSize at begin is ...%ld.\n",rowSize);
+    size_t rowSize = mxGroup->mnLength; //, srcSize = 0;
+    fprintf(stderr,"rowSize at begin is ...%ld.\n",(long)rowSize);
     int *rangeStart =NULL; // The first position for calculation,for example,the A1 in (=MAX(A1:A100))
     int *rangeEnd = NULL; // The last position for calculation,for example, the A100 in (=MAX(A1:A100))
     // The row quantity can be gotten from p2->GetArrayLength()
@@ -236,21 +236,8 @@ bool FormulaGroupInterpreter::interpret()
 //                fprintf(stderr,"After GPU,rRsults[%d] is ...%f\n",i,rResult[i]);
 //            }
 
-// We want to stuff the double data, which in rResult[i] from GPU calculated well, to UI view for users
-            ScAddress aInsertPos = maTopPos;
-            for (sal_Int32 i = 0; i < mxGroup->mnLength; ++i)
-            {
-                aInsertPos.SetRow(mxGroup->mnStart + i);
-                ScFormulaCell* pDestx = mrDoc.GetFormulaCell(aInsertPos);
-
-                SAL_DEBUG(" put value " << rResult[i] << " into formula at " << aInsertPos.Col() << " , " << aInsertPos.Row() );
-                assert(pDestx);
-
-                formula::FormulaTokenRef xResult = new formula::FormulaDoubleToken(rResult[i]);
-                pDestx->SetResultToken(xResult.get());
-                pDestx->ResetDirty();
-                pDestx->SetChanged(true);
-            }
+            // Insert the double data, in rResult[i] back into the document
+            mrDoc.SetFormulaResults(maTopPos, rResult, mxGroup->mnLength);
         }
 
         if(leftData)
commit f4ea3fd5c1ee0e26be8383b0cc209a0d8de3f433
Author: Michael Meeks <michael.meeks at suse.com>
Date:   Wed Jun 26 14:04:25 2013 +0100

    use #if to avoid gotchas, move header to include, make UI conditional.

diff --git a/config_host/config_features.h.in b/config_host/config_features.h.in
index e288eef..fac764e 100644
--- a/config_host/config_features.h.in
+++ b/config_host/config_features.h.in
@@ -59,4 +59,11 @@
 
 #define HAVE_FEATURE_MULTIUSER_ENVIRONMENT 0
 
+/*
+ * Whether we have the OpenCL headers and should compile in any
+ * support for that abstraction.
+ */
+
+#define HAVE_FEATURE_OPENCL 0
+
 #endif
diff --git a/configure.ac b/configure.ac
index 12fe94a..4615ef3 100644
--- a/configure.ac
+++ b/configure.ac
@@ -9854,6 +9854,7 @@ else
                OPENCL_CFLAGS="-I$with_opencl_sdk/include"
                OPENCL_LIBS="-L$with_opencl_sdk/lib/x86 -lOpenCL"
                AC_MSG_RESULT([found at path $with_opencl_sdk])
+               AC_DEFINE(HAVE_FEATURE_OPENCL)
        else
                AC_MSG_ERROR([no headers found found at $with_opencl_sdk/include ])
        fi
diff --git a/sc/source/core/opencl/openclwrapper.hxx b/sc/source/core/inc/openclwrapper.hxx
similarity index 99%
rename from sc/source/core/opencl/openclwrapper.hxx
rename to sc/source/core/inc/openclwrapper.hxx
index d3b5354..62006d1 100644
--- a/sc/source/core/opencl/openclwrapper.hxx
+++ b/sc/source/core/inc/openclwrapper.hxx
@@ -10,6 +10,8 @@
 #ifndef _OPENCL_WRAPPER_H_
 #define _OPENCL_WRAPPER_H_
 
+#include <config_features.h>
+
 #include <CL/cl.h>
 
 #define MaxTextExtent  4096
@@ -22,7 +24,7 @@
 #define strcasecmp strcmp
 #endif
 #endif
-#define ENABLE_OPENCL //dbg
+
 typedef struct _KernelEnv {
     cl_context context;
     cl_command_queue commandQueue;
diff --git a/sc/source/core/tool/formulagroup.cxx b/sc/source/core/tool/formulagroup.cxx
index 43977ed..2262cc5 100644
--- a/sc/source/core/tool/formulagroup.cxx
+++ b/sc/source/core/tool/formulagroup.cxx
@@ -7,6 +7,7 @@
  * file, You can obtain one at http://mozilla.org/MPL/2.0/.
  */
 
+#include <config_features.h>
 #include "formulagroup.hxx"
 #include "document.hxx"
 #include "formulacell.hxx"
@@ -15,8 +16,8 @@
 #include "interpre.hxx"
 #include "formula/vectortoken.hxx"
 
-#ifdef ENABLE_OPENCL
-#include "openclwrapper.hxx"
+#if HAVE_FEATURE_OPENCL
+#  include "openclwrapper.hxx"
 #endif
 
 namespace sc {
@@ -40,7 +41,7 @@ TimeValue aTimeBefore, aTimeAfter;
 
 bool FormulaGroupInterpreter::interpret()
 {
-#ifdef ENABLE_OPENCL //dbg
+#if HAVE_FEATURE_OPENCL
     size_t rowSize = mxGroup->mnLength, srcSize = 0;
     fprintf(stderr,"rowSize at begin is ...%ld.\n",rowSize);
     int *rangeStart =NULL; // The first position for calculation,for example,the A1 in (=MAX(A1:A100))
@@ -100,7 +101,7 @@ bool FormulaGroupInterpreter::interpret()
                         nRowEnd += i;
                     size_t nRowSize = nRowEnd - nRowStart + 1;
                     ScMatrixRef pMat(new ScMatrix(nColSize, nRowSize, 0.0));
-#ifdef ENABLE_OPENCL
+#if HAVE_FEATURE_OPENCL
                     //srcSize = rowSize+nRowSize-rowSize%nRowSize;//align as nRowSize
                     //srcData = (double *)calloc(srcSize,sizeof(double));
                     rangeStart[i] = nRowStart;//record the start position
@@ -109,7 +110,7 @@ bool FormulaGroupInterpreter::interpret()
                     for (size_t nCol = 0; nCol < nColSize; ++nCol)
                     {
                         const double* pArray = rArrays[nCol];
-#ifdef ENABLE_OPENCL
+#if HAVE_FEATURE_OPENCL
                         //printf("pArray is %p.\n",pArray);
                         if( NULL==pArray )
                         {
@@ -146,7 +147,7 @@ bool FormulaGroupInterpreter::interpret()
         if (!pDest)
             return false;
 
-#ifdef ENABLE_OPENCL
+#if HAVE_FEATURE_OPENCL
         const formula::FormulaToken *pCur = aCode2.First();
         aCode2.Reset();
         while( ( pCur = aCode2.Next() ) != NULL )
@@ -187,7 +188,7 @@ bool FormulaGroupInterpreter::interpret()
         }
     } // for loop end (mxGroup->mnLength)
     // For GPU calculation
-#ifdef ENABLE_OPENCL //dbg: Using "export SC_GPU=1" to open if{} in terminal
+#if HAVE_FEATURE_OPENCL //dbg: Using "export SC_GPU=1" to open if{} in terminal
     if(getenv("SC_GPU"))
     {
             fprintf(stderr,"ggGPU flow...\n\n");
@@ -236,17 +237,20 @@ bool FormulaGroupInterpreter::interpret()
 //            }
 
 // We want to stuff the double data, which in rResult[i] from GPU calculated well, to UI view for users
-                for (sal_Int32 i = 0; i < mxGroup->mnLength; ++i)
-                    {
-                    ScFormulaCell* pDestx = mrDoc.GetFormulaCell(aTmpPos);
-                    if (!pDestx)
-                        return false;
-                    formula::FormulaTokenRef xResult = new formula::FormulaDoubleToken(rResult[i]);
-                    pDestx->SetResultToken(xResult.get());
-                    pDestx->ResetDirty();
-                    pDestx->SetChanged(true);
-                    aTmpPos.SetRow(mxGroup->mnStart + i + 1);
-                 }
+            ScAddress aInsertPos = maTopPos;
+            for (sal_Int32 i = 0; i < mxGroup->mnLength; ++i)
+            {
+                aInsertPos.SetRow(mxGroup->mnStart + i);
+                ScFormulaCell* pDestx = mrDoc.GetFormulaCell(aInsertPos);
+
+                SAL_DEBUG(" put value " << rResult[i] << " into formula at " << aInsertPos.Col() << " , " << aInsertPos.Row() );
+                assert(pDestx);
+
+                formula::FormulaTokenRef xResult = new formula::FormulaDoubleToken(rResult[i]);
+                pDestx->SetResultToken(xResult.get());
+                pDestx->ResetDirty();
+                pDestx->SetChanged(true);
+            }
         }
 
         if(leftData)
diff --git a/sc/source/ui/app/scmod.cxx b/sc/source/ui/app/scmod.cxx
index 673b79c..4edfcf0 100644
--- a/sc/source/ui/app/scmod.cxx
+++ b/sc/source/ui/app/scmod.cxx
@@ -17,6 +17,8 @@
  *   the License at http://www.apache.org/licenses/LICENSE-2.0 .
  */
 
+#include <config_features.h>
+
 #include <com/sun/star/ui/dialogs/XSLTFilterDialog.hpp>
 #include <comphelper/processfactory.hxx>
 
@@ -101,8 +103,8 @@
 #include "scabstdlg.hxx"
 #include "formula/errorcodes.hxx"
 
-#ifdef ENABLE_OPENCL
-#include "openclwrapper.hxx"
+#if HAVE_FEATURE_OPENCL
+#  include "openclwrapper.hxx"
 #endif
 
 #define SC_IDLE_MIN     150
@@ -152,7 +154,7 @@ ScModule::ScModule( SfxObjectFactory* pFact ) :
     mbIsInSharedDocLoading( false ),
     mbIsInSharedDocSaving( false )
 {
-#ifdef ENABLE_OPENCL
+#if HAVE_FEATURE_OPENCL
     OclCalc::InitEnv();
 #endif
     //  im ctor ist der ResManager (DLL-Daten) noch nicht initialisiert!
@@ -188,7 +190,7 @@ ScModule::ScModule( SfxObjectFactory* pFact ) :
 
 ScModule::~ScModule()
 {
-#ifdef ENABLE_OPENCL
+#if HAVE_FEATURE_OPENCL
     OclCalc::ReleaseOpenclRunEnv();
 #endif
     OSL_ENSURE( !pSelTransfer, "Selection Transfer object not deleted" );
diff --git a/sc/source/ui/optdlg/calcoptionsdlg.cxx b/sc/source/ui/optdlg/calcoptionsdlg.cxx
index 86d718a..11ec4cb 100644
--- a/sc/source/ui/optdlg/calcoptionsdlg.cxx
+++ b/sc/source/ui/optdlg/calcoptionsdlg.cxx
@@ -10,6 +10,8 @@
  *
  */
 
+#include <config_features.h>
+
 #include "calcoptionsdlg.hxx"
 #include "sc.hrc"
 #include "scresid.hxx"
@@ -197,7 +199,9 @@ void ScCalcOptionsDialog::FillOptionsList()
     }
 
     pModel->Insert(createBoolItem(maCaptionEmptyStringAsZero,maConfig.mbEmptyStringAsZero));
+#if HAVE_FEATURE_OPENCL
     pModel->Insert(createBoolItem(maCaptionOpenCLEnabled,maConfig.mbOpenCLEnabled));
+#endif
 
     mpLbSettings->SetUpdateMode(true);
 }
commit b4be964fa2eece48e7ac1ae7c2c2690890ee899c
Author: Michael Meeks <michael.meeks at suse.com>
Date:   Wed Jun 26 13:07:47 2013 +0100

    cleanup conditionals.

diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index b06af59..4790510 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -10,10 +10,11 @@
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
+#include "sal/config.h"
 #include "random.hxx"
 #include "openclwrapper.hxx"
 #include "oclkernels.hxx"
-#ifdef WIN32
+#ifdef SAL_WIN32
 #include <Windows.h>
 #endif
 //#define USE_KERNEL_FILE
@@ -21,7 +22,7 @@ using namespace std;
 GPUEnv OpenclDevice::gpuEnv;
 int OpenclDevice::isInited =0;
 
-#ifdef WIN32
+#ifdef SAL_WIN32
 
 #define OPENCL_DLL_NAME "opencllo.dll"
 #define OCLERR -1
@@ -54,8 +55,9 @@ void OpenclDevice::FreeOpenclDll()
 
 int OpenclDevice::InitEnv()
 {
-#ifdef WIN32
-	while(1){
+#ifdef SAL_WIN32
+	while(1)
+    {
 	    if(1==LoadOpencl())
 			break;
 	}
@@ -66,7 +68,7 @@ int OpenclDevice::InitEnv()
 
 int OpenclDevice::ReleaseOpenclRunEnv() {
 	ReleaseOpenclEnv(&gpuEnv);
-#ifdef WIN32
+#ifdef SAL_WIN32
 	FreeOpenclDll();
 #endif
     return 1;
@@ -738,16 +740,18 @@ int OpenclDevice::RegisterKernelWrapper(const char *kernelName,cl_kernel_functio
 			return (1);
 		}
 	}
-		return (0);
+    return (0);
 }
 
 
-void OpenclDevice::SetOpenclState(int state) {
+void OpenclDevice::SetOpenclState(int state)
+{
      //printf("OpenclDevice::setOpenclState...\n");
      isInited = state;
 }
 
-int OpenclDevice::GetOpenclState() {
+int OpenclDevice::GetOpenclState()
+{
     return isInited;
 }
 //ocldbg
@@ -886,6 +890,7 @@ int OclFormulaxDll(void ** usrdata, KernelEnv *env) {
 
     return 0;
 }
+
 double OclCalc::OclProcess(cl_kernel_function function, double *data, formulax type)
 {
 	fprintf(stderr, "\n OpenclDevice, proc...begin\n");
@@ -930,6 +935,7 @@ OclCalc::~OclCalc()
     OpenclDevice::SetOpenclState(0);
     fprintf(stderr,"OclCalc:: opencl end ok.\n");
 }
+
 /////////////////////////////////////////////////////////////////////////////
 int OclCalc::OclHostFormulaMax(double *srcData,int *start,int *end,double *output,int size) {
 	KernelEnv env;
diff --git a/sc/source/core/tool/formulagroup.cxx b/sc/source/core/tool/formulagroup.cxx
index 88c33cf..43977ed 100644
--- a/sc/source/core/tool/formulagroup.cxx
+++ b/sc/source/core/tool/formulagroup.cxx
@@ -173,7 +173,8 @@ bool FormulaGroupInterpreter::interpret()
 //            }
         }
 #endif
-        if(getenv("SC_FORMULAGROUP")&&(!getenv("SC_GPU"))){
+        if(!getenv("SC_GPU"))
+        {
             fprintf(stderr,"ccCPU flow...\n\n");
             ScCompiler aComp(&mrDoc, aTmpPos, aCode2);
             aComp.SetGrammar(mrDoc.GetGrammar());
@@ -186,11 +187,12 @@ bool FormulaGroupInterpreter::interpret()
         }
     } // for loop end (mxGroup->mnLength)
     // For GPU calculation
-#ifdef ENABLE_OPENCL //dbg: Using "export SC_FORMULAGROUP=1;export SC_GPU=1" to open if{} in terminal
-    if(getenv("SC_FORMULAGROUP")&&(getenv("SC_GPU"))){
+#ifdef ENABLE_OPENCL //dbg: Using "export SC_GPU=1" to open if{} in terminal
+    if(getenv("SC_GPU"))
+    {
             fprintf(stderr,"ggGPU flow...\n\n");
             printf(" oclOp is... %d\n",oclOp);
-osl_getSystemTime(&aTimeBefore);//timer
+            osl_getSystemTime(&aTimeBefore); //timer
             static OclCalc ocl_calc;
             switch(oclOp)
             {
@@ -219,14 +221,13 @@ osl_getSystemTime(&aTimeBefore);//timer
                     fprintf(stderr,"No OpenCL function for this calculation.\n");
                     break;
             }
-/////////////////////////////////////////////////////
-osl_getSystemTime(&aTimeAfter);
-double diff = getTimeDiff(aTimeAfter, aTimeBefore);
-//if (diff >= 1.0)
-{
-    fprintf(stderr,"OpenCL,diff...%f.\n",diff);
-
-}
+            /////////////////////////////////////////////////////
+            osl_getSystemTime(&aTimeAfter);
+            double diff = getTimeDiff(aTimeAfter, aTimeBefore);
+            //if (diff >= 1.0)
+            {
+                fprintf(stderr,"OpenCL,diff...%f.\n",diff);
+            }
 /////////////////////////////////////////////////////
 
 //rResult[i];
@@ -261,10 +262,10 @@ double diff = getTimeDiff(aTimeAfter, aTimeBefore);
         if(srcData)
             free(srcData);
 
-if(getenv("SC_GPUSAMPLE")){
-    //fprintf(stderr,"FormulaGroupInterpreter::interpret(),iniflag...%d\n",ocl_calc.GetOpenclState());
-    //ocl_calc.OclTest();//opencl test sample for debug
-}
+        if(getenv("SC_GPUSAMPLE")){
+            //fprintf(stderr,"FormulaGroupInterpreter::interpret(),iniflag...%d\n",ocl_calc.GetOpenclState());
+            //ocl_calc.OclTest();//opencl test sample for debug
+        }
 #endif
 
     return true;
commit 4ca21e22bf90159228a5fa001dba9514cb0916c6
Author: Michael Meeks <michael.meeks at suse.com>
Date:   Wed Jun 26 12:52:53 2013 +0100

    avoid srand / time / rand calls.

diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index bcb9cd8..b06af59 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -900,8 +900,6 @@ double OclCalc::OclProcess(cl_kernel_function function, double *data, formulax t
 double OclCalc::OclTest() {
     double data[NUM];
 
-    srand((unsigned int) time(NULL));
-
     for (int i = 0; i < NUM; i++) {
         data[i] = sc::rng::uniform();
         fprintf(stderr, "%f\t", data[i]);
@@ -912,10 +910,9 @@ double OclCalc::OclTest() {
 
 double OclCalc::OclTestDll() {
     double data[NUM];
-    srand((unsigned int) time(NULL));
 
     for (int i = 0; i < NUM; i++) {
-        data[i] = rand() / (RAND_MAX + 1.0);
+        data[i] = sc::rng::uniform();
         fprintf(stderr, "%f\t", data[i]);
     }
     OclProcess(&OclFormulaxDll, data, AVG);
commit 6ae1656b44f3217eddfe9c047e116c1071cd734a
Author: Jing Xian <jingxian at multicorewareinc.com>
Date:   Wed Jun 26 12:19:51 2013 +0100

    more work on formula interpretation.

diff --git a/sc/source/core/opencl/oclkernels.hxx b/sc/source/core/opencl/oclkernels.hxx
index f9db447..3269f3a 100644
--- a/sc/source/core/opencl/oclkernels.hxx
+++ b/sc/source/core/opencl/oclkernels.hxx
@@ -6,17 +6,15 @@
  * License, v. 2.0. If a copy of the MPL was not distributed with this
  * file, You can obtain one at http://mozilla.org/MPL/2.0/.
  */
-
 #ifndef _OCL_KERNEL_H_
 #define _OCL_KERNEL_H_
 
 #ifndef USE_EXTERNAL_KERNEL
 #define KERNEL( ... )# __VA_ARGS__
 
+
 /////////////////////////////////////////////
 const char *kernel_src = KERNEL(
-//Add kernel here
-//use \n ... \n to define macro
 __kernel void hello(__global uint *buffer)
 
 {
@@ -27,83 +25,134 @@ buffer[idx]=idx;
 }
 
 __kernel void oclformula(__global float *data,
-                       const uint type)
+					   const uint type)
 {
-    const unsigned int i = get_global_id(0);
-
-    switch (type)
-    {
-        case 0:          //MAX
-        {
-            //printf("%i ?%f>%f\n",i,data[2*i],data[2*i+1]);
-            if(data[2*i]>data[2*i+1])
-                data[i] = data[2*i];
-            else
-                data[i] = data[2*i+1];
-            break;
-        }
-        case 1:          //MIN
-        {
-            //printf("%d ?%d<%d\n",i,data[2*i],data[2*i+1]);
-            if(data[2*i]<data[2*i+1])
-                data[i] = data[2*i];
-            else
-                data[i] = data[2*i+1];
-            break;
-        }
-        case 2:          //SUM
-        case 3:          //AVG
-        {
-            //printf("%d %d+%d\n",i,data[2*i],data[2*i+1]);
-            data[i] = data[2*i] + data[2*i+1];
-            break;
-        }
-        default:
-            break;
-
-    }
+	const unsigned int i = get_global_id(0);
+
+	switch (type)
+	{
+		case 0:          //MAX
+		{
+			//printf("%i ?%f>%f\n",i,data[2*i],data[2*i+1]);
+			if(data[2*i]>data[2*i+1])
+				data[i] = data[2*i];
+			else
+				data[i] = data[2*i+1];
+			break;
+		}
+		case 1:          //MIN
+		{
+			//printf("%d ?%d<%d\n",i,data[2*i],data[2*i+1]);
+			if(data[2*i]<data[2*i+1])
+				data[i] = data[2*i];
+			else
+				data[i] = data[2*i+1];
+			break;
+		}
+		case 2:          //SUM
+		case 3:          //AVG
+		{
+			//printf("%d %d+%d\n",i,data[2*i],data[2*i+1]);
+			data[i] = data[2*i] + data[2*i+1];
+			break;
+		}
+		default:
+			break;
+
+	}
+}
 
+
+__kernel void oclSignedAdd(__global float *ltData,__global float *rtData,__global float *otData)
+{
+	const unsigned int id = get_global_id(0);
+	otData[id] = ltData[id] + rtData[id];
 }
 
-__kernel void oclFormulaMin(__global float *data,
-                            const uint type)
+
+__kernel void oclSignedSub(__global float *ltData,__global float *rtData,__global float *otData)
 {
+	const unsigned int id = get_global_id(0);
+	otData[id] = ltData[id] - rtData[id];
 
 }
 
-__kernel void oclFormulaMax(__global float *data,
-                            const uint type)
+__kernel void oclSignedMul(__global float *ltData,__global float *rtData,__global float *otData)
 {
+	int id = get_global_id(0);
+	otData[id] =ltData[id] * rtData[id];
+}
+
+__kernel void oclSignedDiv(__global float *ltData,__global float *rtData,__global float *otData)
+{
+	const unsigned int id = get_global_id(0);
+	otData[id] = ltData[id] / rtData[id];
+}
+
+__kernel void oclFormulaMin(__global float *input,__global int *start,__global int *end,__global float *output)
+{
+	const unsigned int id = get_global_id(0);
+	int i=0;
+	unsigned int startFlag = start[id];
+	unsigned int endFlag = end[id];
+	float min = input[startFlag];
+	for(i=startFlag;i<=endFlag;i++)
+	{
+		if(input[i]<min)
+			min = input[i];
+	}
+	output[id] = min;
+
+}
+
+__kernel void oclFormulaMax(__global float *input,__global int *start,__global int *end,__global float *output)
+{
+	const unsigned int id = get_global_id(0);
+	int i=0;
+	unsigned int startFlag = start[id];
+	unsigned int endFlag = end[id];
+	float max = input[startFlag];
+	for(i=startFlag;i<=endFlag;i++)
+	{
+		if(input[i]>max)
+			max = input[i];
+	}
+	output[id] = max;
 
 }
 
 __kernel void oclFormulaSum(__global float *data,
-                            const uint type)
+					   const uint type)
 {
 
 }
 
 __kernel void oclFormulaCount(__global float *data,
-                              const uint type)
+					   const uint type)
 {
 
 }
 
-__kernel void oclFormulaAverage(__global float *data,
-                                const uint type)
+__kernel void oclFormulaAverage(__global float *input,__global int *start,__global int *end,__global float *output)
 {
+	const unsigned int id = get_global_id(0);
+	int i=0;
+	float sum=0;
+	for(i = start[id];i<=end[id];i++)
+		sum += input[i];
+	output[id] = sum / (end[id]-start[id]+1);
 
 }
 
 
 __kernel void oclFormulaSumproduct(__global float *data,
-                                   const uint type)
+					   const uint type)
 {
 
 }
 
 __kernel void oclFormulaMinverse(__global float *data,
-                                 const uint type)
+					   const uint type)
 {
 
 }
@@ -112,5 +161,4 @@ __kernel void oclFormulaMinverse(__global float *data,
 
 #endif // USE_EXTERNAL_KERNEL
 #endif //_OCL_KERNEL_H_
-
 /* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 6c3935e..bcb9cd8 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -10,57 +10,140 @@
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
-
 #include "random.hxx"
 #include "openclwrapper.hxx"
 #include "oclkernels.hxx"
+#ifdef WIN32
+#include <Windows.h>
+#endif
+//#define USE_KERNEL_FILE
+using namespace std;
+GPUEnv OpenclDevice::gpuEnv;
+int OpenclDevice::isInited =0;
 
+#ifdef WIN32
 
-inline int OpenclDevice::add_kernel_cfg(int kCount, const char *kName) {
-    strcpy(gpu_env.kernel_names[kCount], kName);
-    gpu_env.kernel_count++;
-    return 0;
+#define OPENCL_DLL_NAME "opencllo.dll"
+#define OCLERR -1
+#define OCLSUCCESS 1
+HINSTANCE HOpenclDll = NULL;
+    void *OpenclDll = NULL;
+
+int OpenclDevice::LoadOpencl()
+{
+	//fprintf(stderr, " LoadOpenclDllxx... \n");
+	OpenclDll = static_cast<HINSTANCE>(HOpenclDll);
+	OpenclDll = LoadLibrary(OPENCL_DLL_NAME);
+	if (!static_cast<HINSTANCE>(OpenclDll))
+	{
+		fprintf(stderr, " Load opencllo.dll failed! \n");
+		FreeLibrary(static_cast<HINSTANCE>(OpenclDll));
+		return OCLERR;
+	}
+	fprintf(stderr, " Load opencllo.dll successfully!\n");
+	return OCLSUCCESS;
 }
 
-int OpenclDevice::regist_opencl_kernel() {
-    if (!gpu_env.isUserCreated) {
-        memset(&gpu_env, 0, sizeof(gpu_env));
-    }
+void OpenclDevice::FreeOpenclDll()
+{
+	fprintf(stderr, " Free opencllo.dll ... \n");
+	if(!static_cast<HINSTANCE>(OpenclDll))
+		FreeLibrary(static_cast<HINSTANCE>(OpenclDll));
+}
+#endif
 
-    gpu_env.file_count = 0; //argc;
-    gpu_env.kernel_count = 0UL;
-
-    add_kernel_cfg(0, (const char*) "hello");
-    add_kernel_cfg(1, (const char*) "oclformula");
-    add_kernel_cfg(2, (const char*) "oclFormulaMin");
-    add_kernel_cfg(3, (const char*) "oclFormulaMax");
-    add_kernel_cfg(4, (const char*) "oclFormulaSum");
-    add_kernel_cfg(5, (const char*) "oclFormulaCount");
-    add_kernel_cfg(6, (const char*) "oclFormulaAverage");
-    add_kernel_cfg(7, (const char*) "oclFormulaSumproduct");
-    add_kernel_cfg(8, (const char*) "oclFormulaMinverse");
+int OpenclDevice::InitEnv()
+{
+#ifdef WIN32
+	while(1){
+	    if(1==LoadOpencl())
+			break;
+	}
+#endif
+	InitOpenclRunEnv(0,NULL);
+	return 1;
+}
+
+int OpenclDevice::ReleaseOpenclRunEnv() {
+	ReleaseOpenclEnv(&gpuEnv);
+#ifdef WIN32
+	FreeOpenclDll();
+#endif
+    return 1;
+}
+///////////////////////////////////////////////////////
+///////////////////////////////////////////////////////
+inline int OpenclDevice::AddKernelConfig(int kCount, const char *kName) {
+    strcpy(gpuEnv.kernelNames[kCount], kName);
+    gpuEnv.kernelCount++;
     return 0;
 }
-OpenclDevice::OpenclDevice() :
-        isInited(0) {
 
+int OpenclDevice::RegistOpenclKernel() {
+    if (!gpuEnv.isUserCreated) {
+        memset(&gpuEnv, 0, sizeof(gpuEnv));
+    }
+
+    gpuEnv.fileCount = 0; //argc;
+    gpuEnv.kernelCount = 0UL;
+
+    AddKernelConfig(0, (const char*) "hello");
+    AddKernelConfig(1, (const char*) "oclformula");
+    AddKernelConfig(2, (const char*) "oclFormulaMin");
+    AddKernelConfig(3, (const char*) "oclFormulaMax");
+    AddKernelConfig(4, (const char*) "oclFormulaSum");
+    AddKernelConfig(5, (const char*) "oclFormulaCount");
+    AddKernelConfig(6, (const char*) "oclFormulaAverage");
+    AddKernelConfig(7, (const char*) "oclFormulaSumproduct");
+    AddKernelConfig(8, (const char*) "oclFormulaMinverse");
+
+    AddKernelConfig(9,  (const char*) "oclSignedAdd");
+    AddKernelConfig(10, (const char*) "oclSignedSub");
+    AddKernelConfig(11, (const char*) "oclSignedMul");
+    AddKernelConfig(12, (const char*) "oclSignedDiv");
+	return 0;
+}
+OpenclDevice::OpenclDevice(){
+	//InitEnv();
 }
 
 OpenclDevice::~OpenclDevice() {
+	//ReleaseOpenclRunEnv();
+}
 
+int OpenclDevice::CheckKernelName(KernelEnv *envInfo,const char *kernelName){
+    //printf("CheckKernelName,total count of kernels...%d\n", gpuEnv.kernelCount);
+    int kCount;
+    for(kCount=0; kCount < gpuEnv.kernelCount; kCount++) {
+        if(strcasecmp(kernelName, gpuEnv.kernelNames[kCount]) == 0) {
+	    printf("match  %s kernel right\n",kernelName);
+	    break;
+        }
+    }
+    envInfo->context      = gpuEnv.context;
+    envInfo->commandQueue = gpuEnv.commandQueue;
+    envInfo->program      = gpuEnv.programs[0];
+    envInfo->kernel       = gpuEnv.kernels[kCount];
+    strcpy(envInfo->kernelName, kernelName);
+    if (envInfo == (KernelEnv *) NULL)
+    {
+        printf("get err func and env\n");
+        return 0;
+    }
+    return 1;
 }
-#ifdef USE_KERNEL_FILE
-int OpenclDevice::convert_to_string(const char *filename, char **source) {
+
+int OpenclDevice::ConvertToString(const char *filename, char **source) {
     int file_size;
     size_t result;
     FILE *file = NULL;
-
     file_size = 0;
     result = 0;
     file = fopen(filename, "rb+");
-    printf("open kernel file %s.\n", filename);
+    printf("open kernel file %s.\n",filename);
 
     if (file != NULL) {
+		printf("Open ok!\n");
         fseek(file, 0, SEEK_END);
 
         file_size = ftell(file);
@@ -82,68 +165,41 @@ int OpenclDevice::convert_to_string(const char *filename, char **source) {
     printf("open kernel file failed.\n");
     return (0);
 }
-#endif
-int OpenclDevice::binary_generated(cl_context context,
-        const char * cl_file_name, FILE ** fhandle) {
-    unsigned int i = 0;
-    cl_int status;
-
-    size_t numDevices;
-
-    cl_device_id *devices;
-
-    char *str = NULL;
-
-    FILE *fd = NULL;
-
-    status = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES,
-            sizeof(numDevices), &numDevices, NULL);
-
-    CHECK_OPENCL(status)
-
-    devices = (cl_device_id*) malloc(sizeof(cl_device_id) * numDevices);
-
-    if (devices == NULL) {
-        return 0;
-    }
-
-    /* grab the handles to all of the devices in the context. */
-    status = clGetContextInfo(context, CL_CONTEXT_DEVICES,
-            sizeof(cl_device_id) * numDevices, devices, NULL);
-
-    status = 0;
-    /* dump out each binary into its own separate file. */
-    for (i = 0; i < numDevices; i++) {
-        char fileName[256] = { 0 }, cl_name[128] = { 0 };
-
-        if (devices[i] != 0) {
-            char deviceName[1024];
-            status = clGetDeviceInfo(devices[i], CL_DEVICE_NAME,
-                    sizeof(deviceName), deviceName, NULL);
-            CHECK_OPENCL(status)
-            str = (char*) strstr(cl_file_name, (char*) ".cl");
-            memcpy(cl_name, cl_file_name, str - cl_file_name);
-            cl_name[str - cl_file_name] = '\0';
-            sprintf(fileName, "./%s-%s.bin", cl_name, deviceName);
-            fd = fopen(fileName, "rb");
-            status = (fd != NULL) ? 1 : 0;
-        }
 
-    }
-
-    if (devices != NULL) {
-        free(devices);
-        devices = NULL;
-    }
-
-    if (fd != NULL) {
-        *fhandle = fd;
-    }
+int OpenclDevice::BinaryGenerated(const char * clFileName, FILE ** fhandle) {
+    unsigned int i = 0;
+	cl_int status;
+	char *str = NULL;
+	FILE *fd = NULL;
+	cl_uint numDevices=0;
+	status = clGetDeviceIDs(gpuEnv.platform, // platform
+							CL_DEVICE_TYPE_GPU, // device_type
+							0, // num_entries
+							NULL, // devices
+							&numDevices);
+	for (i = 0; i <numDevices; i++) {
+		char fileName[256] = { 0 }, cl_name[128] = { 0 };
+		if (gpuEnv.devices[i] != 0) {
+			char deviceName[1024];
+			status = clGetDeviceInfo(gpuEnv.devices[i], CL_DEVICE_NAME,sizeof(deviceName), deviceName, NULL);
+			CHECK_OPENCL(status);
+			str = (char*) strstr(clFileName, (char*) ".cl");
+			memcpy(cl_name, clFileName, str - clFileName);
+			cl_name[str - clFileName] = '\0';
+			sprintf(fileName, "./%s-%s.bin", cl_name, deviceName);
+			fd = fopen(fileName, "rb");
+			status = (fd != NULL) ? 1 : 0;
+			}
+		}
+		if (fd != NULL) {
+			*fhandle = fd;
+			}
+
+		return status;
 
-    return status;
 }
 
-int OpenclDevice::write_binary_to_file(const char* fileName, const char* birary,
+int OpenclDevice::WriteBinaryToFile(const char* fileName, const char* birary,
         size_t numBytes) {
     FILE *output = NULL;
     output = fopen(fileName, "wb");
@@ -155,11 +211,12 @@ int OpenclDevice::write_binary_to_file(const char* fileName, const char* birary,
     fclose(output);
 
     return 1;
+
 }
 
-int OpenclDevice::generat_bin_from_kernel_source(cl_program program,
-        const char * cl_file_name) {
-    unsigned int i = 0;
+int OpenclDevice::GeneratBinFromKernelSource(cl_program program,
+        const char * clFileName) {
+     unsigned int i = 0;
     cl_int status;
     size_t *binarySizes, numDevices;
     cl_device_id *devices;
@@ -216,12 +273,12 @@ int OpenclDevice::generat_bin_from_kernel_source(cl_program program,
                     sizeof(deviceName), deviceName, NULL);
             CHECK_OPENCL(status)
 
-            str = (char*) strstr(cl_file_name, (char*) ".cl");
-            memcpy(cl_name, cl_file_name, str - cl_file_name);
-            cl_name[str - cl_file_name] = '\0';
+            str = (char*) strstr(clFileName, (char*) ".cl");
+            memcpy(cl_name, clFileName, str - clFileName);
+            cl_name[str - clFileName] = '\0';
             sprintf(fileName, "./%s-%s.bin", cl_name, deviceName);
 
-            if (!write_binary_to_file(fileName, binaries[i], binarySizes[i])) {
+            if (!WriteBinaryToFile(fileName, binaries[i], binarySizes[i])) {
                 printf("opencl-wrapper: write binary[%s] failds\n", fileName);
                 return 0;
             } //else
@@ -254,164 +311,36 @@ int OpenclDevice::generat_bin_from_kernel_source(cl_program program,
     return 1;
 }
 
-int OpenclDevice::init_opencl_attr(OpenCLEnv * env) {
-    if (gpu_env.isUserCreated) {
+int OpenclDevice::InitOpenclAttr(OpenCLEnv * env) {
+    if (gpuEnv.isUserCreated) {
         return 1;
     }
 
-    gpu_env.context = env->context;
-    gpu_env.platform = env->platform;
-    gpu_env.dev = env->devices;
-    gpu_env.command_queue = env->command_queue;
+    gpuEnv.context = env->context;
+    gpuEnv.platform = env->platform;
+    gpuEnv.dev = env->devices;
+    gpuEnv.commandQueue = env->commandQueue;
 
-    gpu_env.isUserCreated = 1;
+    gpuEnv.isUserCreated = 1;
 
     return 0;
 }
 
-int OpenclDevice::create_kernel(char * kernelname, KernelEnv * env) {
+int OpenclDevice::CreateKernel(char * kernelname, KernelEnv * env) {
     int status;
 
-    env->kernel = clCreateKernel(gpu_env.programs[0], kernelname, &status);
-    env->context = gpu_env.context;
-    env->command_queue = gpu_env.command_queue;
+    env->kernel = clCreateKernel(gpuEnv.programs[0], kernelname, &status);
+    env->context = gpuEnv.context;
+    env->commandQueue = gpuEnv.commandQueue;
     return status != CL_SUCCESS ? 1 : 0;
 }
 
-int OpenclDevice::release_kernel(KernelEnv * env) {
+int OpenclDevice::ReleaseKernel(KernelEnv * env) {
     int status = clReleaseKernel(env->kernel);
     return status != CL_SUCCESS ? 1 : 0;
 }
 
-int OpenclDevice::init_opencl_env(GPUEnv *gpu_info) {
-    size_t length;
-    cl_int status;
-    cl_uint numPlatforms, numDevices;
-    cl_platform_id *platforms;
-    cl_context_properties cps[3];
-    char platformName[100];
-    unsigned int i;
-
-    /*
-     * Have a look at the available platforms.
-     */
-    if (!gpu_info->isUserCreated) {
-        status = clGetPlatformIDs(0, NULL, &numPlatforms);
-        if (status != CL_SUCCESS) {
-            return (1);
-        }
-        gpu_info->platform = NULL;
-        ;
-        if (0 < numPlatforms) {
-            platforms = (cl_platform_id*) malloc(
-                    numPlatforms * sizeof(cl_platform_id));
-            if (platforms == (cl_platform_id*) NULL) {
-                return (1);
-            }
-            status = clGetPlatformIDs(numPlatforms, platforms, NULL);
-
-            if (status != CL_SUCCESS) {
-                return (1);
-            }
-
-            for (i = 0; i < numPlatforms; i++) {
-                status = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR,
-                        sizeof(platformName), platformName, NULL);
-
-                if (status != CL_SUCCESS) {
-                    return (1);
-                }
-                gpu_info->platform = platforms[i];
-
-                //if (!strcmp(platformName, "Intel(R) Coporation"))
-                //if( !strcmp( platformName, "Advanced Micro Devices, Inc." ))
-                {
-                    gpu_info->platform = platforms[i];
-
-                    status = clGetDeviceIDs(gpu_info->platform /* platform */,
-                            CL_DEVICE_TYPE_GPU /* device_type */,
-                            0 /* num_entries */, NULL /* devices */,
-                            &numDevices);
-
-                    if (status != CL_SUCCESS) {
-                        return (1);
-                    }
-
-                    if (numDevices) {
-                        break;
-                    }
-                }
-            }
-            free(platforms);
-        }
-        if (NULL == gpu_info->platform) {
-            return (1);
-        }
-
-        /*
-         * Use available platform.
-         */
-        cps[0] = CL_CONTEXT_PLATFORM;
-        cps[1] = (cl_context_properties) gpu_info->platform;
-        cps[2] = 0;
-        /* Check for GPU. */
-        gpu_info->dType = CL_DEVICE_TYPE_GPU;
-        gpu_info->context = clCreateContextFromType(cps, gpu_info->dType, NULL,
-                NULL, &status);
-
-        if ((gpu_info->context == (cl_context) NULL)
-                || (status != CL_SUCCESS)) {
-            gpu_info->dType = CL_DEVICE_TYPE_CPU;
-            gpu_info->context = clCreateContextFromType(cps, gpu_info->dType,
-                    NULL, NULL, &status);
-        }
-        if ((gpu_info->context == (cl_context) NULL)
-                || (status != CL_SUCCESS)) {
-            gpu_info->dType = CL_DEVICE_TYPE_DEFAULT;
-            gpu_info->context = clCreateContextFromType(cps, gpu_info->dType,
-                    NULL, NULL, &status);
-        }
-        if ((gpu_info->context == (cl_context) NULL)
-                || (status != CL_SUCCESS)) {
-            return (1);
-        }
-        /* Detect OpenCL devices. */
-        /* First, get the size of device list data */
-        status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES, 0,
-                NULL, &length);
-
-        if ((status != CL_SUCCESS) || (length == 0)) {
-            return (1);
-        }
-        /* Now allocate memory for device list based on the size we got earlier */
-        gpu_info->devices = (cl_device_id*) malloc(length);
-        if (gpu_info->devices == (cl_device_id*) NULL) {
-            return (1);
-        }
-        /* Now, get the device list data */
-        status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES, length,
-                gpu_info->devices, NULL);
-
-        if (status != CL_SUCCESS) {
-            return (1);
-        }
-
-        /* Create OpenCL command queue. */
-        gpu_info->command_queue = clCreateCommandQueue(gpu_info->context,
-                gpu_info->devices[0], 0, &status);
-
-        if (status != CL_SUCCESS) {
-            return (1);
-        }
-    }
-
-    status = clGetCommandQueueInfo(gpu_info->command_queue,
-            CL_QUEUE_THREAD_HANDLE_AMD, 0, NULL, NULL);
-
-    return 0;
-}
-
-int OpenclDevice::release_opencl_env(GPUEnv *gpu_info) {
+int OpenclDevice::ReleaseOpenclEnv(GPUEnv *gpuInfo) {
     int i = 0;
     int status = 0;
 
@@ -419,60 +348,44 @@ int OpenclDevice::release_opencl_env(GPUEnv *gpu_info) {
         return 1;
     }
 
-    for (i = 0; i < gpu_env.file_count; i++) {
-        if (gpu_env.programs[i]) {
-            status = clReleaseProgram(gpu_env.programs[i]);
+    for (i = 0; i < gpuEnv.fileCount; i++) {
+        if (gpuEnv.programs[i]) {
+            status = clReleaseProgram(gpuEnv.programs[i]);
             CHECK_OPENCL(status)
-            gpu_env.programs[i] = NULL;
+            gpuEnv.programs[i] = NULL;
         }
     }
-    if (gpu_env.command_queue) {
-        clReleaseCommandQueue(gpu_env.command_queue);
-        gpu_env.command_queue = NULL;
+    if (gpuEnv.commandQueue) {
+        clReleaseCommandQueue(gpuEnv.commandQueue);
+        gpuEnv.commandQueue = NULL;
     }
-    if (gpu_env.context) {
-        clReleaseContext(gpu_env.context);
-        gpu_env.context = NULL;
+    if (gpuEnv.context) {
+        clReleaseContext(gpuEnv.context);
+        gpuEnv.context = NULL;
     }
     isInited = 0;
-    gpu_info->isUserCreated = 0;
-    free(gpu_info->devices);
+    gpuInfo->isUserCreated = 0;
+    free(gpuInfo->devices);
     return 1;
 }
 
-int OpenclDevice::run_kernel_wrapper(cl_kernel_function function,
-        char * kernel_name, void **usrdata) {
-    printf("oclwrapper:run_kernel_wrapper...\n");
-    if (register_kernel_wrapper(kernel_name, function) != 1) {
+int OpenclDevice::RunKernelWrapper(cl_kernel_function function,
+        const char * kernelName, void **usrdata) {
+    printf("oclwrapper:RunKernel_wrapper...\n");
+    if (RegisterKernelWrapper(kernelName, function) != 1) {
         fprintf(stderr,
-                "Error:run_kernel_wrapper:register_kernel_wrapper fail!\n");
+                "Error:RunKernel_wrapper:RegisterKernelWrapper fail!\n");
         return -1;
     }
-    return (run_kernel(kernel_name, usrdata));
-}
-
-int OpenclDevice::register_kernel_wrapper(const char *kernel_name,
-        cl_kernel_function function) {
-    int i;
-    printf("oclwrapper:register_kernel_wrapper...%d\n", gpu_env.kernel_count);
-    for (i = 0; i < gpu_env.kernel_count; i++) {
-        //printf("oclwrapper:register_kernel_wrapper kname...%s\n", kernel_name);
-        //printf("oclwrapper:register_kernel_wrapper kname...%s\n", gpu_env.kernel_names[i]);
-        if (strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
-            //printf("oclwrapper:register_kernel_wrapper if()...\n");
-            gpu_env.kernel_functions[i] = function;
-            return (1);
-        }
-    }
-    return (0);
+    return (RunKernel(kernelName, usrdata));
 }
 
-int OpenclDevice::cached_of_kerner_prg(const GPUEnv *gpu_env_cached,
-        const char * cl_file_name) {
-    int i;
-    for (i = 0; i < gpu_env_cached->file_count; i++) {
-        if (strcasecmp(gpu_env_cached->kernelSrcFile[i], cl_file_name) == 0) {
-            if (gpu_env_cached->programs[i] != NULL) {
+int OpenclDevice::CachedOfKernerPrg(const GPUEnv *gpuEnvCached,
+        const char * clFileName) {
+  int i;
+    for (i = 0; i < gpuEnvCached->fileCount; i++) {
+        if (strcasecmp(gpuEnvCached->kernelSrcFile[i], clFileName) == 0) {
+            if (gpuEnvCached->programs[i] != NULL) {
                 return (1);
             }
         }
@@ -481,37 +394,30 @@ int OpenclDevice::cached_of_kerner_prg(const GPUEnv *gpu_env_cached,
     return (0);
 }
 
-int OpenclDevice::compile_kernel_file(GPUEnv *gpu_info, const char *build_option) {
+int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) {
     cl_int status;
-
     size_t length;
-
     char *buildLog = NULL, *binary;
-
     const char *source;
     size_t source_size[1];
-
     int b_error, binary_status, binaryExisted, idx;
-
     size_t numDevices;
-
     cl_device_id *devices;
-
     FILE *fd, *fd1;
     const char* filename = "kernel.cl";
-    if (cached_of_kerner_prg(gpu_info, filename) == 1) {
+	fprintf(stderr, "CompileKernelFile ... \n");
+    if (CachedOfKernerPrg(gpuInfo, filename) == 1) {
         return (1);
     }
 
-    idx = gpu_info->file_count;
+    idx = gpuInfo->fileCount;
 
     source = kernel_src;
 
     source_size[0] = strlen(source);
-
     binaryExisted = 0;
-    if ((binaryExisted = binary_generated(gpu_info->context, filename, &fd)) == 1) {
-        status = clGetContextInfo(gpu_info->context, CL_CONTEXT_NUM_DEVICES,
+    if ((binaryExisted = BinaryGenerated(filename, &fd)) == 1) {
+        status = clGetContextInfo(gpuInfo->context, CL_CONTEXT_NUM_DEVICES,
                 sizeof(numDevices), &numDevices, NULL);
         CHECK_OPENCL(status)
 
@@ -543,11 +449,11 @@ int OpenclDevice::compile_kernel_file(GPUEnv *gpu_info, const char *build_option
         fclose(fd);
         fd = NULL;
         // grab the handles to all of the devices in the context.
-        status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES,
+        status = clGetContextInfo(gpuInfo->context, CL_CONTEXT_DEVICES,
                 sizeof(cl_device_id) * numDevices, devices, NULL);
         CHECK_OPENCL(status)
 
-        gpu_info->programs[idx] = clCreateProgramWithBinary(gpu_info->context,
+        gpuInfo->programs[idx] = clCreateProgramWithBinary(gpuInfo->context,
                 numDevices, devices, &length, (const unsigned char**) &binary,
                 &binary_status, &status);
         CHECK_OPENCL(status)
@@ -556,40 +462,37 @@ int OpenclDevice::compile_kernel_file(GPUEnv *gpu_info, const char *build_option
         free(devices);
         devices = NULL;
     } else {
-
         // create a CL program using the kernel source
-        gpu_info->programs[idx] = clCreateProgramWithSource(gpu_info->context,
+        gpuEnv.programs[idx] = clCreateProgramWithSource(gpuEnv.context,
                 1, &source, source_size, &status);
-        CHECK_OPENCL(status)
-
-        printf("clCreateProgramWithSource.\n");
+        CHECK_OPENCL(status);
     }
 
-    if (gpu_info->programs[idx] == (cl_program) NULL) {
+    if (gpuInfo->programs[idx] == (cl_program) NULL) {
         return (0);
     }
 
     //char options[512];
     // create a cl program executable for all the devices specified
-    if (!gpu_info->isUserCreated) {
-        status = clBuildProgram(gpu_info->programs[idx], 1, gpu_info->devices,
-                build_option, NULL, NULL);
+    if (!gpuInfo->isUserCreated) {
+        status = clBuildProgram(gpuInfo->programs[idx], 1, gpuInfo->devices,
+                buildOption, NULL, NULL);
         CHECK_OPENCL(status)
     } else {
-        status = clBuildProgram(gpu_info->programs[idx], 1, &(gpu_info->dev),
-                build_option, NULL, NULL);
+        status = clBuildProgram(gpuInfo->programs[idx], 1, &(gpuInfo->dev),
+                buildOption, NULL, NULL);
         CHECK_OPENCL(status)
     }
     printf("BuildProgram.\n");
 
     if (status != CL_SUCCESS) {
-        if (!gpu_info->isUserCreated) {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, 0, NULL,
+        if (!gpuInfo->isUserCreated) {
+            status = clGetProgramBuildInfo(gpuInfo->programs[idx],
+                    gpuInfo->devices[0], CL_PROGRAM_BUILD_LOG, 0, NULL,
                     &length);
         } else {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
+            status = clGetProgramBuildInfo(gpuInfo->programs[idx],
+                    gpuInfo->dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
         }
         if (status != CL_SUCCESS) {
             printf("opencl create build log fail\n");
@@ -599,13 +502,13 @@ int OpenclDevice::compile_kernel_file(GPUEnv *gpu_info, const char *build_option
         if (buildLog == (char*) NULL) {
             return (0);
         }
-        if (!gpu_info->isUserCreated) {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, length,
+        if (!gpuInfo->isUserCreated) {
+            status = clGetProgramBuildInfo(gpuInfo->programs[idx],
+                    gpuInfo->devices[0], CL_PROGRAM_BUILD_LOG, length,
                     buildLog, &length);
         } else {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->dev, CL_PROGRAM_BUILD_LOG, length, buildLog,
+            status = clGetProgramBuildInfo(gpuInfo->programs[idx],
+                    gpuInfo->dev, CL_PROGRAM_BUILD_LOG, length, buildLog,
                     &length);
         }
 
@@ -619,199 +522,35 @@ int OpenclDevice::compile_kernel_file(GPUEnv *gpu_info, const char *build_option
         return (0);
     }
 
-    strcpy(gpu_env.kernelSrcFile[idx], filename);
+    strcpy(gpuEnv.kernelSrcFile[idx], filename);
 
     if (binaryExisted == 0)
-        generat_bin_from_kernel_source(gpu_env.programs[idx], filename);
+        GeneratBinFromKernelSource(gpuEnv.programs[idx], filename);
 
-    gpu_info->file_count += 1;
+    gpuInfo->fileCount += 1;
 
     return (1);
-}
-
-int OpenclDevice::compile_kernel_file(const char *filename, GPUEnv *gpu_info,
-        const char *build_option) {
-    cl_int status;
-
-    size_t length;
-
-#ifdef USE_KERNEL_FILE
-    char
-    *source_str;
-#endif
-    char *buildLog = NULL, *binary;
-
-    const char *source;
-    size_t source_size[1];
-
-    int b_error, binary_status, binaryExisted, idx;
-
-    size_t numDevices;
-
-    cl_device_id *devices;
-
-    FILE *fd, *fd1;
-
-    if (cached_of_kerner_prg(gpu_info, filename) == 1) {
-        return (1);
-    }
-
-    idx = gpu_info->file_count;
-#ifdef USE_KERNEL_FILE
-    status = convert_to_string( filename, &source_str, gpu_info, idx );
-
-    if( status == 0 )
-    {
-        printf("convert_to_string failed.\n");
-        return(0);
-    }
-    source = source_str;
-#else
-
-    source = kernel_src;
-#endif
-    source_size[0] = strlen(source);
-
-    binaryExisted = 0;
-    if ((binaryExisted = binary_generated(gpu_info->context, filename, &fd))
-            == 1) {
-        status = clGetContextInfo(gpu_info->context, CL_CONTEXT_NUM_DEVICES,
-                sizeof(numDevices), &numDevices, NULL);
-        CHECK_OPENCL(status)
-
-        devices = (cl_device_id*) malloc(sizeof(cl_device_id) * numDevices);
-        if (devices == NULL) {
-            return 0;
-        }
-
-        b_error = 0;
-        length = 0;
-        b_error |= fseek(fd, 0, SEEK_END) < 0;
-        b_error |= (length = ftell(fd)) <= 0;
-        b_error |= fseek(fd, 0, SEEK_SET) < 0;
-        if (b_error) {
-            return 0;
-        }
-
-        binary = (char*) malloc(length + 2);
-        if (!binary) {
-            return 0;
-        }
-
-        memset(binary, 0, length + 2);
-        b_error |= fread(binary, 1, length, fd) != length;
-        if (binary[length - 1] != '\n') {
-            binary[length++] = '\n';
-        }
-
-        fclose(fd);
-        fd = NULL;
-        /* grab the handles to all of the devices in the context. */
-        status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES,
-                sizeof(cl_device_id) * numDevices, devices, NULL);
-        CHECK_OPENCL(status)
-
-        gpu_info->programs[idx] = clCreateProgramWithBinary(gpu_info->context,
-                numDevices, devices, &length, (const unsigned char**) &binary,
-                &binary_status, &status);
-        CHECK_OPENCL(status)
-
-        free(binary);
-        free(devices);
-        devices = NULL;
-    } else {
-
-        // create a CL program using the kernel source
-        gpu_info->programs[idx] = clCreateProgramWithSource(gpu_info->context,
-                1, &source, source_size, &status);
-        CHECK_OPENCL(status)
-#ifdef USE_KERNEL_FILE
-        free((char*)source);
-#endif
-        printf("clCreateProgramWithSource.\n");
-    }
-
-    if (gpu_info->programs[idx] == (cl_program) NULL) {
-        return (0);
-    }
-
-    //char options[512];
-    // create a cl program executable for all the devices specified
-    if (!gpu_info->isUserCreated) {
-        status = clBuildProgram(gpu_info->programs[idx], 1, gpu_info->devices,
-                build_option, NULL, NULL);
-        CHECK_OPENCL(status)
-    } else {
-        status = clBuildProgram(gpu_info->programs[idx], 1, &(gpu_info->dev),
-                build_option, NULL, NULL);
-        CHECK_OPENCL(status)
-    }
-    printf("BuildProgram.\n");
-
-    if (status != CL_SUCCESS) {
-        if (!gpu_info->isUserCreated) {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, 0, NULL,
-                    &length);
-        } else {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
-        }
-        if (status != CL_SUCCESS) {
-            printf("opencl create build log fail\n");
-            return (0);
-        }
-        buildLog = (char*) malloc(length);
-        if (buildLog == (char*) NULL) {
-            return (0);
-        }
-        if (!gpu_info->isUserCreated) {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, length,
-                    buildLog, &length);
-        } else {
-            status = clGetProgramBuildInfo(gpu_info->programs[idx],
-                    gpu_info->dev, CL_PROGRAM_BUILD_LOG, length, buildLog,
-                    &length);
-        }
-
-        fd1 = fopen("kernel-build.log", "w+");
-        if (fd1 != NULL) {
-            fwrite(buildLog, sizeof(char), length, fd1);
-            fclose(fd1);
-        }
 
-        free(buildLog);
-        return (0);
-    }
 
-    strcpy(gpu_env.kernelSrcFile[idx], filename);
-
-    if (binaryExisted == 0)
-        generat_bin_from_kernel_source(gpu_env.programs[idx], filename);
-
-    gpu_info->file_count += 1;
-
-    return (1);
 }
-
-int OpenclDevice::get_kernel_env_and_func(const char *kernel_name,
+int OpenclDevice::GetKernelEnvAndFunc(const char *kernelName,
         KernelEnv *env, cl_kernel_function *function) {
     int i; //,program_idx ;
-    for (i = 0; i < gpu_env.kernel_count; i++) {
-        if (strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
-            env->context = gpu_env.context;
-            env->command_queue = gpu_env.command_queue;
-            env->program = gpu_env.programs[0];
-            env->kernel = gpu_env.kernels[i];
-            *function = gpu_env.kernel_functions[i];
+    printf("----------------OpenclDevice::GetKernelEnvAndFunc\n");
+    for (i = 0; i < gpuEnv.kernelCount; i++) {
+        if (strcasecmp(kernelName, gpuEnv.kernelNames[i]) == 0) {
+            env->context = gpuEnv.context;
+            env->commandQueue = gpuEnv.commandQueue;
+            env->program = gpuEnv.programs[0];
+            env->kernel = gpuEnv.kernels[i];
+            *function = gpuEnv.kernelFunctions[i];
             return (1);
         }
     }
     return (0);
 }
 
-int OpenclDevice::run_kernel(const char *kernel_name, void **userdata) {
+int OpenclDevice::RunKernel(const char *kernelName, void **userdata) {
     KernelEnv env;
 
     cl_kernel_function function;
@@ -819,8 +558,8 @@ int OpenclDevice::run_kernel(const char *kernel_name, void **userdata) {
     int status;
 
     memset(&env, 0, sizeof(KernelEnv));
-    status = get_kernel_env_and_func(kernel_name, &env, &function);
-    strcpy(env.kernel_name, kernel_name);
+    status = GetKernelEnvAndFunc(kernelName, &env, &function);
+    strcpy(env.kernelName, kernelName);
     if (status == 1) {
         if (&env == (KernelEnv *) NULL
                 || &function == (cl_kernel_function *) NULL) {
@@ -830,11 +569,9 @@ int OpenclDevice::run_kernel(const char *kernel_name, void **userdata) {
     }
     return (0);
 }
-
-int OpenclDevice::init_opencl_run_env(int argc, const char *build_option_kernelfiles)
+int OpenclDevice::InitOpenclRunEnv(int argc, const char *buildOptionKernelfiles)
 {
     int status = 0;
-
     if (MAX_CLKERNEL_NUM <= 0) {
         return 1;
     }
@@ -843,82 +580,177 @@ int OpenclDevice::init_opencl_run_env(int argc, const char *build_option_kernelf
     }
 
     if (!isInited) {
-        printf("regist_opencl_kernel start.\n");
-        regist_opencl_kernel();
+        RegistOpenclKernel();
         //initialize devices, context, comand_queue
-        status = init_opencl_env(&gpu_env);
+        status = InitOpenclRunEnv(&gpuEnv);
         if (status) {
             printf("init_opencl_env failed.\n");
             return (1);
         }
         printf("init_opencl_env successed.\n");
-        //initialize program, kernel_name, kernel_count
-        status = compile_kernel_file( &gpu_env, build_option_kernelfiles);
-        if (status == 0 || gpu_env.kernel_count == 0) {
-            printf("compile_kernel_file failed.\n");
+        //initialize program, kernelName, kernelCount
+        status = CompileKernelFile( &gpuEnv, buildOptionKernelfiles);
+        if (status == 0 || gpuEnv.kernelCount == 0) {
+            printf("CompileKernelFile failed.\n");
             return (1);
         }
-        printf("compile_kernel_file successed.\n");
+        printf("CompileKernelFile successed.\n");
         isInited = 1;
     }
-
     return (0);
 }
 
-int OpenclDevice::init_opencl_run_env(int argc, const char *argv_kernelfiles[],
-        const char *build_option_kernelfiles) {
-    int status = 0;
+int OpenclDevice::InitOpenclRunEnv(GPUEnv *gpuInfo)
+{
+    size_t length;
+    cl_int status;
+    cl_uint numPlatforms, numDevices;
+    cl_platform_id *platforms;
+    cl_context_properties cps[3];
+    char platformName[100];
+    unsigned int i;
 
-    if (MAX_CLKERNEL_NUM <= 0) {
-        return 1;
-    }
-    if ((argc > MAX_CLFILE_NUM) || (argc < 0)) {
-        return 1;
-    }
+    // Have a look at the available platforms.
 
-    if (!isInited) {
-        printf("regist_opencl_kernel start.\n");
-        regist_opencl_kernel();
-        //initialize devices, context, comand_queue
-        status = init_opencl_env(&gpu_env);
-        if (status) {
-            printf("init_opencl_env failed.\n");
+    if (!gpuInfo->isUserCreated) {
+        status = clGetPlatformIDs(0, NULL, &numPlatforms);
+        if (status != CL_SUCCESS) {
             return (1);
         }
-        printf("init_opencl_env successed.\n");
-        //initialize program, kernel_name, kernel_count
-        status = compile_kernel_file(argv_kernelfiles[0], &gpu_env,
-                build_option_kernelfiles);
-        if (status == 0 || gpu_env.kernel_count == 0) {
-            printf("compile_kernel_file failed.\n");
+        gpuInfo->platform = NULL;
+
+        if (0 < numPlatforms) {
+            platforms = (cl_platform_id*) malloc(
+                    numPlatforms * sizeof(cl_platform_id));
+            if (platforms == (cl_platform_id*) NULL) {
+                return (1);
+            }
+            status = clGetPlatformIDs(numPlatforms, platforms, NULL);
+
+            if (status != CL_SUCCESS) {
+                return (1);
+            }
+
+            for (i = 0; i < numPlatforms; i++) {
+                status = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR,
+                        sizeof(platformName), platformName, NULL);
+
+                if (status != CL_SUCCESS) {
+                    return (1);
+                }
+                gpuInfo->platform = platforms[i];
+
+                //if (!strcmp(platformName, "Intel(R) Coporation"))
+                //if( !strcmp( platformName, "Advanced Micro Devices, Inc." ))
+                {
+                    gpuInfo->platform = platforms[i];
+
+                    status = clGetDeviceIDs(gpuInfo->platform, // platform
+												CL_DEVICE_TYPE_GPU, // device_type
+												0, // num_entries
+												NULL, // devices
+												&numDevices);
+
+                    if (status != CL_SUCCESS) {
+                        return (1);
+                    }
+
+                    if (numDevices) {
+                        break;
+                    }
+                }
+            }
+            free(platforms);
+        }
+        if (NULL == gpuInfo->platform) {
+            return (1);
+        }
+
+        // Use available platform.
+
+        cps[0] = CL_CONTEXT_PLATFORM;
+        cps[1] = (cl_context_properties) gpuInfo->platform;
+        cps[2] = 0;
+        // Check for GPU.
+        gpuInfo->dType = CL_DEVICE_TYPE_GPU;
+        gpuInfo->context = clCreateContextFromType(cps, gpuInfo->dType, NULL,
+                NULL, &status);
+
+        if ((gpuInfo->context == (cl_context) NULL)
+                || (status != CL_SUCCESS)) {
+            gpuInfo->dType = CL_DEVICE_TYPE_CPU;
+            gpuInfo->context = clCreateContextFromType(cps, gpuInfo->dType,
+                    NULL, NULL, &status);
+        }
+        if ((gpuInfo->context == (cl_context) NULL)
+                || (status != CL_SUCCESS)) {
+            gpuInfo->dType = CL_DEVICE_TYPE_DEFAULT;
+            gpuInfo->context = clCreateContextFromType(cps, gpuInfo->dType,
+                    NULL, NULL, &status);
+        }
+        if ((gpuInfo->context == (cl_context) NULL)
+                || (status != CL_SUCCESS)) {
+            return (1);
+        }
+        // Detect OpenCL devices.
+        // First, get the size of device list data
+        status = clGetContextInfo(gpuInfo->context, CL_CONTEXT_DEVICES, 0,
+                NULL, &length);
+        if ((status != CL_SUCCESS) || (length == 0)) {
+            return (1);
+        }
+        // Now allocate memory for device list based on the size we got earlier
+        gpuInfo->devices = (cl_device_id*) malloc(length);
+        if (gpuInfo->devices == (cl_device_id*) NULL) {
+            return (1);
+        }
+        // Now, get the device list data
+        status = clGetContextInfo(gpuInfo->context, CL_CONTEXT_DEVICES, length,
+                gpuInfo->devices, NULL);
+        if (status != CL_SUCCESS) {
+            return (1);
+        }
+
+        // Create OpenCL command queue.
+        gpuInfo->commandQueue = clCreateCommandQueue(gpuInfo->context,
+                gpuInfo->devices[0], 0, &status);
+
+        if (status != CL_SUCCESS) {
             return (1);
         }
-        printf("compile_kernel_file successed.\n");
-        isInited = 1;
     }
 
-    return (0);
-}
+    status = clGetCommandQueueInfo(gpuInfo->commandQueue,
+            CL_QUEUE_THREAD_HANDLE_AMD, 0, NULL, NULL);
+
+    return 0;
 
-int OpenclDevice::release_opencl_run_env() {
-    return release_opencl_env(&gpu_env);
 }
+int OpenclDevice::RegisterKernelWrapper(const char *kernelName,cl_kernel_function function)
+{
+	int i;
+	printf("oclwrapper:RegisterKernelWrapper...%d\n", gpuEnv.kernelCount);
+	for (i = 0; i < gpuEnv.kernelCount; i++)
+	{
+		if (strcasecmp(kernelName, gpuEnv.kernelNames[i]) == 0)
+		{
+			gpuEnv.kernelFunctions[i] = function;
+			return (1);
+		}
+	}
+		return (0);
+}
+
 
-void OpenclDevice::setOpenclState(int state) {
+void OpenclDevice::SetOpenclState(int state) {
+     //printf("OpenclDevice::setOpenclState...\n");
      isInited = state;
 }
 
-int OpenclDevice::getOpenclState() {
+int OpenclDevice::GetOpenclState() {
     return isInited;
 }
 //ocldbg
-int OclFormulaMin(void ** usrdata, KernelEnv *env) { return 0; }
-int OclFormulaMax(void ** usrdata, KernelEnv *env) { return 0; }
-int OclFormulaSum(void ** usrdata, KernelEnv *env) { return 0; }
-int OclFormulaCount(void ** usrdata, KernelEnv *env) { return 0; }
-int OclFormulaAverage(void ** usrdata, KernelEnv *env) { return 0; }
-int OclFormulaSumproduct(void ** usrdata, KernelEnv *env) { return 0; }
-int OclFormulaMinverse(void ** usrdata, KernelEnv *env) { return 0; }
 
 int OclFormulax(void ** usrdata, KernelEnv *env) {
     fprintf(stderr, "In OpenclDevice,...Formula_proc\n");
@@ -958,16 +790,16 @@ int OclFormulax(void ** usrdata, KernelEnv *env) {
 
     while (global_work_size[0] != 1) {
         global_work_size[0] = global_work_size[0] / 2;
-        status = clEnqueueNDRangeKernel(env->command_queue, env->kernel, 1,
+        status = clEnqueueNDRangeKernel(env->commandQueue, env->kernel, 1,
                 NULL, global_work_size, NULL, 0, NULL, NULL);
         CHECK_OPENCL(status)
 
     }
     //fprintf(stderr, "\nIn OpenclDevice,...before clEnqueueReadBuffer\n");
-    status = clEnqueueReadBuffer(env->command_queue, formula_data, CL_FALSE, 0,
+    status = clEnqueueReadBuffer(env->commandQueue, formula_data, CL_FALSE, 0,
             sizeof(float), (void *) &tdata, 0, NULL, NULL);
     CHECK_OPENCL(status)
-    status = clFinish(env->command_queue);
+    status = clFinish(env->commandQueue);
     CHECK_OPENCL(status)
 
     //PPAStopCpuEvent(ppa_proc);
@@ -986,46 +818,572 @@ int OclFormulax(void ** usrdata, KernelEnv *env) {
 
     return 0;
 }
-double OclCalc::OclProcess(cl_kernel_function function, double *data,
-        formulax type) {
-    fprintf(stderr, "\In OpenclDevice, proc...begin\n");
-    double ret = 0;
 
-    void *usrdata[2];
+int OclFormulaxDll(void ** usrdata, KernelEnv *env) {
+
+    fprintf(stderr, "In OclFormulaxDll...\n");
+    cl_int clStatus;
+    int status;
+    size_t global_work_size[1];
+    float tdata[NUM];
+
+    double *data = (double *) usrdata[0];
+    const formulax type = *((const formulax *) usrdata[1]);
+    double ret = 0.0;
 
-    usrdata[0] = (void *) data;
-    usrdata[1] = (void *) &type;
+    for (int i = 0; i < NUM; i++) {
+        tdata[i] = (float) data[i];
+    }
 
-    run_kernel_wrapper(function, "oclformula", usrdata);
-    //fprintf(stderr, "\In OpenclDevice, proc...after run_kernel_wrapper\n");
-    return ret;
+    env->kernel = clCreateKernel(env->program, "oclformula", &clStatus);
+    //printf("ScInterpreter::IterateParameters...after clCreateKernel.\n");
+    //fprintf(stderr, "\nIn OpenclDevice,...after clCreateKernel\n");
+    int size = NUM;
+
+    cl_mem formula_data = clCreateBuffer(env->context,
+            (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR),
+            size * sizeof(float), (void *) tdata, &clStatus);
+    //fprintf(stderr, "\nIn OpenclDevice,...after clCreateBuffer\n");
+
+    status = clSetKernelArg(env->kernel, 0, sizeof(cl_mem),
+            (void *) &formula_data);
+    CHECK_OPENCL(status)
+    status = clSetKernelArg(env->kernel, 1, sizeof(unsigned int),
+            (void *) &type);
+    CHECK_OPENCL(status)
+
+    global_work_size[0] = size;
+    //fprintf(stderr, "\nIn OpenclDevice,...after global_work_size\n");
+    //PPAStartCpuEvent(ppa_proc);
+
+    while (global_work_size[0] != 1) {
+        global_work_size[0] = global_work_size[0] / 2;
+        status = clEnqueueNDRangeKernel(env->commandQueue, env->kernel, 1,
+                NULL, global_work_size, NULL, 0, NULL, NULL);
+        CHECK_OPENCL(status)
+
+    }
+    //fprintf(stderr, "\nIn OpenclDevice,...before clEnqueueReadBuffer\n");
+    status = clEnqueueReadBuffer(env->commandQueue, formula_data, CL_FALSE, 0,
+            sizeof(float), (void *) &tdata, 0, NULL, NULL);
+    CHECK_OPENCL(status)
+    status = clFinish(env->commandQueue);
+    CHECK_OPENCL(status)
+
+    //PPAStopCpuEvent(ppa_proc);
+    //fprintf(stderr, "\nIn OpenclDevice,...before clReleaseKernel\n");

... etc. - the rest is truncated


More information about the Libreoffice-commits mailing list