[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