[Libreoffice-commits] core.git: Branch 'feature/calc-group-interpreter-2' - 2 commits - sc/inc sc/qa sc/source
Haidong Lian
haidong at multicorewareinc.com
Fri Aug 30 12:46:35 PDT 2013
sc/inc/column.hxx | 1
sc/inc/document.hxx | 15
sc/qa/unit/filters-test.cxx | 9
sc/source/core/data/column3.cxx | 11
sc/source/core/data/documentimport.cxx | 1
sc/source/core/data/formulacell.cxx | 4
sc/source/core/data/table6.cxx | 2
sc/source/core/opencl/formulagroupcl.cxx | 1149 ++++++++++++++++++--------
sc/source/core/opencl/oclkernels.hxx | 239 ++++-
sc/source/core/opencl/openclwrapper.cxx | 1333 ++++++++++++++++++++++++-------
sc/source/core/opencl/openclwrapper.hxx | 182 ++--
sc/source/filter/excel/excform.cxx | 7
sc/source/filter/excel/read.cxx | 2
sc/source/filter/oox/workbookhelper.cxx | 1
14 files changed, 2205 insertions(+), 751 deletions(-)
New commits:
commit 6da9e1ccd8b07ee8270397afda1a6fa404507af4
Author: Haidong Lian <haidong at multicorewareinc.com>
Date: Fri Aug 30 15:35:17 2013 -0400
Patch for milestone1-0829-v4.
1. Add the parser based on RPN;
2. For test sample1 named "ground-water-daily.xls", using the compound formula
to do calculation;
Add the compound kernels:
Formulae include "AVERAGE,MAX and MIN".Compound formulae include "AVERAGE
-(+,*,/)","MAX -(+,*,/)" and "MIN -(+,*,/)";
3. For formulae which do not work in GPU, they'll work in CPU;
4. For compound operators(-,+,*,/), they'll be calculated one by one in GPU as
the sequence of RPN;
5. Add the start and end position to fit for the sliding window;
6. Modify kernels by using vector for AMD GPU.
Change-Id: I6157008575ce89ddd3e7bf552a87812474af4125
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index 915f1d9..a835c46 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -19,6 +19,10 @@
#include "openclwrapper.hxx"
+#define SRCDATASIZE 100
+#define SINGLEARRAYLEN 100
+#define DOUBLEARRAYLEN 100
+#define SVDOUBLELEN 100
namespace sc {
// A single public entry point for a factory function:
@@ -38,30 +42,632 @@ double getTimeDiff(const TimeValue& t1, const TimeValue& t2)
}//dbg-t
TimeValue aTimeBefore, aTimeAfter;
///////////////////////////////////////
+class SourceData
+{
+ const double *mdpSrcPtr;
+ unsigned int mnDataSize;
+ const char *mcpFormulaName;
+ unsigned int mnCol;
+ int eOp;
+public:
+ SourceData( const double *dpData, unsigned int nSize, uint nCol = 1,const char *cpFormulaName = NULL):mdpSrcPtr(dpData),mnDataSize(nSize),mcpFormulaName(cpFormulaName),mnCol(nCol)
+ {
+ }
+ SourceData():mdpSrcPtr(NULL),mnDataSize(0)
+ {
+ }
+ void setSrcPtr( const double *dpTmpDataPtr)
+ {
+ mdpSrcPtr = dpTmpDataPtr;
+ }
+ void setSrcSize( int nSize )
+ {
+ mnDataSize = nSize;
+ }
+ const double * getDouleData()
+ {
+ return mdpSrcPtr;
+ }
+ unsigned int getDataSize()
+ {
+ return mnDataSize;
+ }
+ void print()
+ {
+ for( uint i=0; i<mnDataSize; i++ )
+ printf( " The SourceData is %f and data size is %d\n",mdpSrcPtr[i],mnDataSize );
+ }
+ void printFormula()
+ {
+ printf("--------The formulaname is %s and the eOp is %d---------\n",mcpFormulaName,eOp);
+ }
+ void setFormulaName(const char *cpFormulaName)
+ {
+ this->mcpFormulaName = cpFormulaName;
+ }
+ const char *getFormulaName()
+ {
+ return mcpFormulaName;
+ }
+ void seteOp(int op)
+ {
+ this->eOp = op;
+ }
+ int geteOp()
+ {
+ return eOp;
+ }
+ int getColNum()
+ {
+ return mnCol;
+ }
+
+};
class FormulaGroupInterpreterOpenCL : public FormulaGroupInterpreterSoftware
{
+ SourceData *mSrcDataStack[SRCDATASIZE];
+ unsigned int mnStackPointer,mnDoublePtrCount;
+ uint * mnpOclStartPos;
+ uint * mnpOclEndPos;
+ SingleVectorFormula *mSingleArray[SINGLEARRAYLEN];
+ DoubleVectorFormula *mDoubleArray[DOUBLEARRAYLEN];
+ double mdpSvdouble[SVDOUBLELEN];
+ double *mdpSrcDoublePtr[SVDOUBLELEN];
+ uint mnSingleCount;
+ uint mnDoubleCount;
+ uint mnSvDoubleCount;
+ uint mnOperatorGroup[100];
+ uint mnOperatorCount;
+ char mcHostName[100];
+ uint mnPositonLen;
+ size_t mnRowSize;
public:
FormulaGroupInterpreterOpenCL() :
FormulaGroupInterpreterSoftware()
{
- OclCalc::InitEnv();
+ mnStackPointer = 0;
+ mnpOclEndPos = NULL;
+ mnpOclStartPos = NULL;
+ mnSingleCount = 0;
+ mnDoubleCount = 0;
+ mnSvDoubleCount = 0;
+ mnOperatorCount = 0;
+ mnPositonLen = 0;
+ mnDoublePtrCount = 0;
+ OclCalc::initEnv();
}
virtual ~FormulaGroupInterpreterOpenCL()
{
- OclCalc::ReleaseOpenclRunEnv();
+ OclCalc::releaseOpenclRunEnv();
}
- virtual ScMatrixRef inverseMatrix(const ScMatrix& rMat);
- virtual bool interpret(ScDocument& rDoc, const ScAddress& rTopPos,
- const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode);
+ virtual ScMatrixRef inverseMatrix( const ScMatrix& rMat );
+ virtual bool interpret( ScDocument& rDoc, const ScAddress& rTopPos,
+ const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode );
+ void collectDoublePointers( double *temp )
+ {
+ if( mnDoublePtrCount < SRCDATASIZE )
+ {
+ mdpSrcDoublePtr[mnDoublePtrCount++] = temp;
+ }
+ else
+ {
+ printf( "The mdpSrcDoublePtr is full now.\n" );
+ double *dtmp = NULL;
+ if ( (dtmp = mdpSrcDoublePtr[--mnDoublePtrCount]) != NULL )
+ {
+ free( dtmp );
+ dtmp = NULL;
+ }
+ }
+ }
+
+ void freeDoublePointers()
+ {
+ while( mnDoublePtrCount > 0 )
+ {
+ double *dtmp = NULL;
+ if ( (dtmp = mdpSrcDoublePtr[--mnDoublePtrCount]) != NULL )
+ {
+ free( dtmp );
+ dtmp = NULL;
+ }
+ }
+ }
+
+
+ void srdDataPush( SourceData *temp )
+ {
+ if( mnStackPointer < SRCDATASIZE )
+ {
+ mSrcDataStack[mnStackPointer++] = temp;
+ }
+ else
+ printf( "The stack is full now.\n" );
+ }
+ SourceData *srdDataPop( void )
+ {
+ if( mnStackPointer <= 0 )
+ {
+ printf( "The stack was empty\n" );
+ return NULL;
+ }
+ return mSrcDataStack[--mnStackPointer];
+ }
+ unsigned int getDataSize()
+ {
+ return mnStackPointer;
+ }
+ void printStackInfo()
+ {
+ printf( "/********The stack size is %d*********\\\n",mnStackPointer );
+ for ( int i = mnStackPointer - 1; i >= 0; i-- )
+ mSrcDataStack[i]->print();
+ }
+ bool getPosition(const ScTokenArray& rCode,const ScFormulaCellGroupRef& xGroup,uint nRowSize,uint *&npOclStartPos,uint *&npOclEndPos,uint *nPositonLen);
+ bool chooseFunction(OclCalc &ocl_calc,double *&dpResult);
+ bool isStockHistory();
+ bool isGroundWater();
};
+bool FormulaGroupInterpreterOpenCL::getPosition(const ScTokenArray& rCode,const ScFormulaCellGroupRef& xGroup,uint nRowSize,uint *&npOclStartPos,uint *&npOclEndPos,uint *nPositonLen)
+{
+ uint nColPosition = 0;
+ ScTokenArray * rCodePos = rCode.Clone();
+ static int nCountPosSize = nRowSize;
+ bool isAllocFormulaOclBuf = true;
+ for ( const formula::FormulaToken* p = rCodePos->First(); p; p = rCodePos->Next() )
+ {
+ switch ( p->GetType() )
+ {
+ case formula::svDoubleVectorRef:
+ {
+ nColPosition++;
+ break;
+ }
+ }
+ }
+ int nPositionSize = nColPosition * nRowSize;
+ npOclStartPos = (unsigned int*) malloc( nPositionSize * sizeof(unsigned int) );
+ npOclEndPos = (unsigned int*) malloc( nPositionSize * sizeof(unsigned int) );
+ if ( nCountPosSize < nPositionSize )
+ {
+ nCountPosSize = nPositionSize;
+ isAllocFormulaOclBuf = false;
+ }
+ for ( sal_Int32 i = 0; i < xGroup->mnLength; ++i )
+ {
+ ScTokenArray * rCodeTemp = rCode.Clone();
+ int j = 0;
+ for ( const formula::FormulaToken* p = rCodeTemp->First(); p; p = rCodeTemp->Next() )
+ {
+ switch (p->GetType())
+ {
+ case formula::svDoubleVectorRef:
+ {
+ const formula::DoubleVectorRefToken* p2 = static_cast<const formula::DoubleVectorRefToken*>(p);
+ size_t nRowStart = p2->IsStartFixed() ? 0 : i;
+ size_t nRowEnd = p2->GetRefRowSize() - 1;
+ if (!p2->IsEndFixed())
+ nRowEnd += i;
+ npOclStartPos[j*nRowSize+i] = nRowStart;//record the start position
+ npOclEndPos[j*nRowSize+i] = nRowEnd;//record the end position
+ j++;
+ }
+ }
+ }
+ }
+ *nPositonLen = nPositionSize;
+ //Now the pos array is 0 1 2 3 4 5 0 1 2 3 4 5;
+ return isAllocFormulaOclBuf;
+}
+
+bool FormulaGroupInterpreterOpenCL::isStockHistory()
+{
+ bool isHistory = false;
+ if( (mnOperatorGroup[0]== 224) && (mnOperatorGroup[1]== 227) && (mnOperatorGroup[2]== 41) && (mnOperatorGroup[3]== 43) && (mnOperatorGroup[4]== 41) )
+ {
+ strcpy(mcHostName,"OclOperationColumnN");
+ isHistory = true;
+ }
+ else if( (mnOperatorGroup[0] == 226) && (mnOperatorGroup[1] == 42) )
+ {
+ strcpy(mcHostName,"OclOperationColumnH");
+ isHistory = true;
+ }
+ else if((mnOperatorGroup[0] == 213) && (mnOperatorGroup[1] == 43) && (mnOperatorGroup[2] == 42) )
+ {
+ strcpy(mcHostName,"OclOperationColumnJ");
+ isHistory = true;
+ }
+ return isHistory;
+}
+
+bool FormulaGroupInterpreterOpenCL::isGroundWater()
+{
+ bool GroundWater=false;
+
+ if((mnOperatorGroup[0] == ocAverage && 1 == mnSingleCount )||(mnOperatorGroup[0] == ocMax && 1 == mnSingleCount )||
+ (mnOperatorGroup[0] == ocMin && 1 == mnSingleCount )||(mnOperatorGroup[0] == ocSub && mnSvDoubleCount==1))
+ {
+ GroundWater = true;
+ }
+ return GroundWater;
+}
-ScMatrixRef FormulaGroupInterpreterOpenCL::inverseMatrix(const ScMatrix& rMat)
+bool FormulaGroupInterpreterOpenCL::chooseFunction( OclCalc &ocl_calc, double *&dpResult )
+{
+ const double * dpOclSrcData = NULL;
+ unsigned int nSrcDataSize = 0;
+ const double *dpLeftData = NULL;
+ const double *dpRightData = NULL;
+ if((mnOperatorGroup[0] == ocAverage && 1 == mnSingleCount )||(mnOperatorGroup[0] == ocMax && 1 == mnSingleCount )||
+ (mnOperatorGroup[0] == ocMin && 1 == mnSingleCount )||(mnOperatorGroup[0] == ocSub && mnSvDoubleCount==1))
+ {
+ double delta = 0.0;
+ const double *pArrayToSubtractOneElementFrom;
+ const double *pGroundWaterDataArray;
+ uint nSrcData = 0;
+ if( mnSvDoubleCount!=1 )
+ {
+ pArrayToSubtractOneElementFrom= mSingleArray[0]->mdpInputLeftData;
+ pGroundWaterDataArray= mDoubleArray[0]->mdpInputData;
+ nSrcData = mDoubleArray[0]->mnInputDataSize;
+ }
+ else
+ {
+ pArrayToSubtractOneElementFrom= mSingleArray[0]->mdpInputLeftData;
+ pGroundWaterDataArray=NULL;
+ delta = mdpSvdouble[0];
+ }
+ ocl_calc.oclGroundWaterGroup( mnOperatorGroup,mnOperatorCount,pGroundWaterDataArray,pArrayToSubtractOneElementFrom,nSrcData,mnRowSize,delta,mnpOclStartPos,mnpOclEndPos,dpResult);
+ }
+ else if( isStockHistory() )
+ {
+ return false;
+ }
+ else if(((mnSvDoubleCount==0)&&(mnSingleCount==0)&&(mnDoubleCount==1)) &&
+ ((mnOperatorGroup[0] == ocAverage)||(mnOperatorGroup[0] == ocMax)||(mnOperatorGroup[0] == ocMin)))
+ {
+ if(mnOperatorGroup[0] == ocAverage)
+ strcpy(mcHostName,"oclFormulaAverage");
+ if(mnOperatorGroup[0] == ocMax)
+ strcpy(mcHostName,"oclFormulaMax");
+ if(mnOperatorGroup[0] == ocMin)
+ strcpy(mcHostName,"oclFormulaMin");
+ DoubleVectorFormula * doubleTemp = mDoubleArray[--mnDoubleCount];
+ nSrcDataSize = doubleTemp->mnInputDataSize;
+ dpOclSrcData = doubleTemp->mdpInputData;
+ if ( ocl_calc.getOpenclState())
+ {
+ if ( ocl_calc.gpuEnv.mnKhrFp64Flag==1 || ocl_calc.gpuEnv.mnAmdFp64Flag == 1 )
+ {
+ ocl_calc.createFormulaBuf64Bits( nSrcDataSize, mnRowSize );
+ ocl_calc.mapAndCopy64Bits( dpOclSrcData,mnpOclStartPos,mnpOclEndPos,nSrcDataSize,mnRowSize );
+ ocl_calc.oclHostFormulaStatistics64Bits( mcHostName, dpResult, mnRowSize );
+ }
+ else
+ {
+ ocl_calc.createFormulaBuf32Bits( nSrcDataSize, mnPositonLen );
+ ocl_calc.mapAndCopy32Bits( dpOclSrcData, mnpOclStartPos, mnpOclEndPos, nSrcDataSize, mnRowSize);
+ ocl_calc.oclHostFormulaStatistics32Bits( mcHostName, dpResult, mnRowSize );
+ }
+ }
+ }
+ else if((mnSvDoubleCount==0)&&(mnSingleCount==1)&&(mnDoubleCount==0))
+ {
+ dpLeftData = mSingleArray[0]->mdpInputLeftData;
+ dpRightData = mSingleArray[0]->mdpInputRightData;
+ if(mnOperatorGroup[0] == ocAdd)
+ strcpy(mcHostName,"oclSignedAdd");
+ if(mnOperatorGroup[0] == ocSub)
+ strcpy(mcHostName,"oclSignedSub");
+ if(mnOperatorGroup[0] == ocMul)
+ strcpy(mcHostName,"oclSignedMul");
+ if(mnOperatorGroup[0] == ocDiv)
+ strcpy(mcHostName,"oclSignedDiv");
+ if ( ocl_calc.getOpenclState())
+ {
+ if ( ocl_calc.gpuEnv.mnKhrFp64Flag == 1 || ocl_calc.gpuEnv.mnAmdFp64Flag == 1 )
+ {
+ ocl_calc.createArithmeticOptBuf64Bits( mnRowSize );
+ ocl_calc.mapAndCopy64Bits(dpLeftData,dpRightData,mnRowSize);
+ ocl_calc.oclHostArithmeticOperator64Bits( mcHostName,dpResult,mnRowSize );
+ }
+ else
+ {
+ ocl_calc.createArithmeticOptBuf32Bits( mnRowSize );
+ ocl_calc.mapAndCopy32Bits(dpLeftData,dpRightData,mnRowSize);
+ ocl_calc.oclHostArithmeticOperator32Bits( mcHostName,dpResult,mnRowSize );
+ }
+ }
+ }
+ else if( (mnSingleCount>1) && (mnSvDoubleCount==0) && (mnDoubleCount==0) )
+ {
+ const double* dpArray[100] = {};
+ int j=0;
+ for( uint i = 0; i < mnSingleCount; i++ )
+ {
+ dpArray[j++] = mSingleArray[i]->mdpInputLeftData;
+ if( NULL != mSingleArray[i]->mdpInputRightData )
+ dpArray[j++] = mSingleArray[i]->mdpInputRightData;
+ }
+ double *dpMoreColArithmetic = (double *)malloc( sizeof(double) * j * mnRowSize );
+ if( NULL == dpMoreColArithmetic )
+ {
+ printf("Memory alloc error!\n");
+ return false;
+ }
+ for( uint i = 0; i < j*mnRowSize; i++ )
+ {
+ dpMoreColArithmetic[i] = dpArray[i/mnRowSize][i%mnRowSize];
+ }
+ if ( ocl_calc.getOpenclState())
+ {
+ if ( ocl_calc.gpuEnv.mnKhrFp64Flag == 1 || ocl_calc.gpuEnv.mnAmdFp64Flag == 1 )
+ {
+ ocl_calc.createMoreColArithmeticBuf64Bits( j * mnRowSize, mnOperatorCount );
+ ocl_calc.mapAndCopyMoreColArithmetic64Bits( dpMoreColArithmetic, mnRowSize * j, mnOperatorGroup, mnOperatorCount );
+ ocl_calc.oclMoreColHostArithmeticOperator64Bits( mnRowSize, mnOperatorCount, dpResult,mnRowSize );
+ }
+ else
+ {
+ ocl_calc.createMoreColArithmeticBuf32Bits( j* mnRowSize, mnOperatorCount );
+ ocl_calc.mapAndCopyMoreColArithmetic32Bits(dpMoreColArithmetic, mnRowSize * j, mnOperatorGroup, mnOperatorCount);
+ ocl_calc.oclMoreColHostArithmeticOperator32Bits( mnRowSize, mnOperatorCount, dpResult, mnRowSize );
+ }
+ }
+ }
+ else
+ {
+ return false;
+ }
+ return true;
+}
+
+class agency
+{
+public:
+ double *calculate(int nOclOp,int rowSize,OclCalc &ocl_calc,uint *npOclStartPos,uint *npOclEndPos,FormulaGroupInterpreterOpenCL *formulaInterprt);
+};
+
+double * agency::calculate( int nOclOp,int rowSize,OclCalc &ocl_calc,uint *npOclStartPos,uint *npOclEndPos,FormulaGroupInterpreterOpenCL *formulaInterprt)
+{
+ const double *dpLeftData = NULL;
+ const double *dpRightData = NULL;
+ const double *dpOclSrcData=NULL;
+ if ( ocl_calc.gpuEnv.mnKhrFp64Flag == 1 || ocl_calc.gpuEnv.mnAmdFp64Flag == 1 )
+ {
+ switch( nOclOp )
+ {
+ case ocAdd:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * nDataSize );
+ memset(rResult,0,rowSize);
+ ocl_calc.oclHostArithmeticStash64Bits( "oclSignedAdd",dpLeftData,dpRightData,rResult,temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ break;
+ }
+ case ocSub:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = ( double * )malloc( sizeof(double) * nDataSize );
+ memset( rResult,0,rowSize );
+ ocl_calc.oclHostArithmeticStash64Bits( "oclSignedSub",dpLeftData,dpRightData,rResult,temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData(rResult,nDataSize) );
+ break;
+ }
+ case ocMul:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * nDataSize );
+ memset( rResult,0,rowSize );
+ ocl_calc.oclHostArithmeticStash64Bits( "oclSignedMul",dpLeftData,dpRightData,rResult,temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ break;
+ }
+ case ocDiv:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = ( double * )malloc( sizeof(double) * nDataSize );
+ memset( rResult,0,rowSize );
+ ocl_calc.oclHostArithmeticStash64Bits( "oclSignedDiv",dpLeftData,dpRightData,rResult,temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ break;
+ }
+ case ocMax:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ nDataSize = temp->getDataSize();
+ dpOclSrcData = temp->getDouleData();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * rowSize );
+ memset( rResult,0,rowSize );
+ ocl_calc.oclHostFormulaStash64Bits( "oclFormulaMax",dpOclSrcData,npOclStartPos,npOclEndPos,rResult,nDataSize,rowSize );
+ formulaInterprt->srdDataPush( new SourceData( rResult,rowSize ) );
+ break;
+ }
+ case ocMin:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ nDataSize = temp->getDataSize();
+ dpOclSrcData = temp->getDouleData();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * rowSize );
+ memset( rResult,0,rowSize );
+ ocl_calc.oclHostFormulaStash64Bits( "oclFormulaMin",dpOclSrcData,npOclStartPos,npOclEndPos,rResult,nDataSize,rowSize );
+ formulaInterprt->srdDataPush( new SourceData( rResult,rowSize ) );
+ break;
+ }
+ case ocAverage:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ nDataSize = temp->getDataSize();
+ dpOclSrcData = temp->getDouleData();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * rowSize );
+ memset( rResult,0,rowSize );
+ ocl_calc.oclHostFormulaStash64Bits( "oclFormulaAverage",dpOclSrcData,npOclStartPos,npOclEndPos,rResult,nDataSize,rowSize );
+ formulaInterprt->srdDataPush( new SourceData( rResult,rowSize ) );
+ break;
+ }
+ default:
+ fprintf( stderr,"No OpenCL function for this calculation.\n" );
+ break;
+ }
+ }
+ else
+ {
+ switch( nOclOp )
+ {
+ case ocAdd:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * nDataSize );
+ memset(rResult,0,rowSize);
+ ocl_calc.oclHostArithmeticStash32Bits( "oclSignedAdd", dpLeftData, dpRightData, rResult, temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData(rResult, nDataSize) );
+ break;
+ }
+ case ocSub:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * nDataSize );
+ memset( rResult, 0, rowSize );
+ ocl_calc.oclHostArithmeticStash32Bits( "oclSignedSub", dpLeftData, dpRightData, rResult, temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ break;
+ }
+ case ocMul:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc(sizeof(double) * nDataSize );
+ memset( rResult, 0, rowSize );
+ ocl_calc.oclHostArithmeticStash32Bits( "oclSignedMul", dpLeftData, dpRightData, rResult, temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData( rResult, nDataSize ) );
+ break;
+ }
+ case ocDiv:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ SourceData *temp2 = formulaInterprt->srdDataPop();
+ nDataSize = temp2->getDataSize();
+ dpLeftData = temp2->getDouleData();
+ dpRightData = temp->getDouleData();
+ nDataSize = temp2->getDataSize();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * nDataSize );
+ memset( rResult, 0, rowSize );
+ ocl_calc.oclHostArithmeticStash32Bits( "oclSignedDiv", dpLeftData, dpRightData, rResult, temp->getDataSize() );
+ formulaInterprt->srdDataPush( new SourceData(rResult, nDataSize) );
+ break;
+ }
+ case ocMax:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ nDataSize = temp->getDataSize();
+ dpOclSrcData = temp->getDouleData();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc(sizeof(double) * nDataSize );
+ memset(rResult,0,rowSize);
+ ocl_calc.oclHostFormulaStash32Bits( "oclFormulaMax", dpOclSrcData, npOclStartPos, npOclEndPos, rResult,nDataSize, rowSize );
+ formulaInterprt->srdDataPush( new SourceData( rResult, rowSize ) );
+ break;
+ }
+ case ocMin:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ nDataSize = temp->getDataSize();
+ dpOclSrcData = temp->getDouleData();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * nDataSize );
+ memset( rResult, 0, rowSize );
+ ocl_calc.oclHostFormulaStash32Bits( "oclFormulaMin", dpOclSrcData, npOclStartPos, npOclEndPos, rResult, nDataSize, rowSize );
+ formulaInterprt->srdDataPush( new SourceData( rResult, rowSize) );
+ break;
+ }
+ case ocAverage:
+ {
+ unsigned int nDataSize = 0;
+ SourceData *temp = formulaInterprt->srdDataPop();
+ nDataSize = temp->getDataSize();
+ dpOclSrcData = temp->getDouleData();
+ double *rResult = NULL; // Point to the output data from GPU
+ rResult = (double *)malloc( sizeof(double) * nDataSize );
+ memset( rResult, 0, rowSize);
+ ocl_calc.oclHostFormulaStash32Bits( "oclFormulaAverage", dpOclSrcData, npOclStartPos, npOclEndPos, rResult, nDataSize, rowSize );
+ formulaInterprt->srdDataPush( new SourceData( rResult, rowSize) );
+ break;
+ }
+ default:
+ fprintf(stderr,"No OpenCL function for this calculation.\n");
+ break;
+ }
+ }
+ return NULL;
+}
+
+ScMatrixRef FormulaGroupInterpreterOpenCL::inverseMatrix( const ScMatrix& rMat )
{
SCSIZE nC, nR;
- rMat.GetDimensions(nC, nR);
- if (nC != nR || nC == 0)
+ rMat.GetDimensions( nC, nR );
+ if ( nC != nR || nC == 0 )
// Input matrix must be square. Return an empty matrix on failure and
// the caller will calculate it via CPU.
return ScMatrixRef();
@@ -70,29 +676,30 @@ ScMatrixRef FormulaGroupInterpreterOpenCL::inverseMatrix(const ScMatrix& rMat)
// the last, chained together in a single array.
std::vector<double> aDoubles;
rMat.GetDoubleArray(aDoubles);
+
float * fpOclMatrixSrc = NULL;
float * fpOclMatrixDst = NULL;
double * dpOclMatrixSrc = NULL;
double * dpOclMatrixDst = NULL;
uint nMatrixSize = nC * nR;
static OclCalc aOclCalc;
- if ( aOclCalc.GetOpenclState() )
+ if ( aOclCalc.getOpenclState() )
{
if ( aOclCalc.gpuEnv.mnKhrFp64Flag == 1 || aOclCalc.gpuEnv.mnAmdFp64Flag == 1 )
{
- aOclCalc.CreateBuffer64Bits( dpOclMatrixSrc, dpOclMatrixDst, nMatrixSize );
+ aOclCalc.createBuffer64Bits( dpOclMatrixSrc, dpOclMatrixDst, nMatrixSize );
for ( uint i = 0; i < nC; i++ )
for ( uint j = 0; j < nR; j++ )
dpOclMatrixSrc[i*nC+j] = aDoubles[j*nR+i];
- aOclCalc.OclHostMatrixInverse64Bits( "oclFormulaMtxInv", dpOclMatrixSrc, dpOclMatrixDst,aDoubles, nR );
+ aOclCalc.oclHostMatrixInverse64Bits( "oclFormulaMtxInv", dpOclMatrixSrc, dpOclMatrixDst,aDoubles, nR );
}
else
{
- aOclCalc.CreateBuffer32Bits( fpOclMatrixSrc, fpOclMatrixDst, nMatrixSize );
+ aOclCalc.createBuffer32Bits( fpOclMatrixSrc, fpOclMatrixDst, nMatrixSize );
for ( uint i = 0; i < nC; i++ )
for ( uint j = 0; j < nR; j++ )
fpOclMatrixSrc[i*nC+j] = (float) aDoubles[j*nR+i];
- aOclCalc.OclHostMatrixInverse32Bits( "oclFormulaMtxInv", fpOclMatrixSrc, fpOclMatrixDst, aDoubles, nR );
+ aOclCalc.oclHostMatrixInverse32Bits( "oclFormulaMtxInv", fpOclMatrixSrc, fpOclMatrixDst, aDoubles, nR );
}
}
@@ -111,372 +718,242 @@ ScMatrixRef FormulaGroupInterpreterOpenCL::inverseMatrix(const ScMatrix& rMat)
#else
// Another way is to put the values one column at a time.
const double* p = &aDoubles[0];
- for (SCSIZE i = 0; i < nC; ++i)
+ for( SCSIZE i = 0; i < nC; ++i )
{
- xInv->PutDouble(p, nR, i, 0);
+ xInv->PutDouble( p, nR, i, 0 );
p += nR;
}
#endif
return xInv;
}
-
-bool FormulaGroupInterpreterOpenCL::interpret(ScDocument& rDoc, const ScAddress& rTopPos,
- const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode)
+bool FormulaGroupInterpreterOpenCL::interpret( ScDocument& rDoc, const ScAddress& rTopPos,
+ const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode )
{
- generateRPNCode(rDoc, rTopPos, rCode);
-
- size_t rowSize = xGroup->mnLength;
- fprintf(stderr,"rowSize at begin is ...%ld.\n",(long)rowSize);
+ generateRPNCode( rDoc, rTopPos, rCode );
+ mnRowSize = xGroup->mnLength;
+ fprintf( stderr,"mnRowSize at begin is ...%ld.\n",(long)mnRowSize );
// The row quantity can be gotten from p2->GetArrayLength()
- uint nCount1 = 0, nCount2 = 0, nCount3 = 0;
int nOclOp = 0;
- double *rResult = NULL; // Point to the output data from GPU
- rResult = (double *)malloc(sizeof(double) * rowSize*2);// For 2 columns(B,C)
- if(NULL==rResult)
- {
- printf("malloc err\n");
- return false;
- }
- memset(rResult,0,rowSize);
- float * fpOclSrcData = NULL; // Point to the input data from CPU
- double * dpOclSrcData = NULL;
- uint * npOclStartPos = NULL; // The first position for calculation,for example,the A1 in (=MAX(A1:A100))
- uint * npOclEndPos = NULL; // The last position for calculation,for example, the A100 in (=MAX(A1:A100))
- float * fpLeftData = NULL; // Left input for binary operator(+,-,*,/),for example,(=leftData+rightData)
- float * fpRightData = NULL; // Right input for binary operator(+,-,*,/),for example,(=leftData/rightData)
- // The rightData can't be zero for "/"
- double * dpLeftData = NULL;
- double * dpRightData = NULL;
-
- float * fpSaveData=NULL; //It is a temp pointer point the preparing memory;
- float * fpSumProMergeLfData = NULL; //It merge the more col to one col is the left operator
- float * fpSumProMergeRtData = NULL; //It merge the more col to one col is the right operator
- double * dpSaveData=NULL;
- double * dpSumProMergeLfData = NULL;
- double * dpSumProMergeRtData = NULL;
- uint * npSumSize=NULL; //It is a array to save the matix sizt(col *row)
- int nSumproductSize=0; //It is the merge array size
- bool aIsAlloc=false; //It is a flag to judge the fpSumProMergeLfData existed
- unsigned int nCountMatix=0; //It is a count to save the calculate times
+ const double * dpOclSrcData = NULL;
+ const double * dpBinaryData = NULL;
static OclCalc ocl_calc;
- bool isSumProduct=false;
- if(ocl_calc.GetOpenclState())
+ unsigned int nSrcDataSize = 0;
+
+ const double *dpResult = NULL;
+ double *pResult = (double *)malloc(sizeof(double) * mnRowSize);
+ double *dpSvDouble = NULL;
+ bool isSample = false;
+
+ mnSingleCount = 0;
+ mnDoubleCount = 0;
+ mnSvDoubleCount = 0;
+ mnOperatorCount = 0;
+ mnPositonLen = 0;
+ if ( ocl_calc.getOpenclState() )
{
- // Don't know how large the size will be applied previously, so create them as the rowSize or 65536
- // Don't know which formulae will be used previously, so create buffers for different formulae used probably
- if(ocl_calc.gpuEnv.mnKhrFp64Flag==1 || ocl_calc.gpuEnv.mnAmdFp64Flag==1)
+ getPosition(rCode,xGroup,mnRowSize,mnpOclStartPos,mnpOclEndPos,&mnPositonLen);
+ const formula::FormulaToken* p = rCode.FirstRPN();
+
+ bool isSingle = false;
+ int nCountNum=0;
+ do
{
- ocl_calc.CreateBuffer64Bits(dpOclSrcData,npOclStartPos,npOclEndPos,rowSize);
- ocl_calc.CreateBuffer64Bits(dpLeftData,dpRightData,rowSize);
- }
- else
+ if ( ocPush != p->GetOpCode())
+ {
+ nOclOp = p->GetOpCode();
+ mnOperatorGroup[mnOperatorCount++] = nOclOp;
+ }
+ else if( ocPush == p->GetOpCode() && formula::svSingleVectorRef == p->GetType() )
+ {
+ mnSingleCount++;
+ }
+ if ( ocPush == p->GetOpCode() && formula::svDouble == p->GetType() )
+ {
+ mnSvDoubleCount++;
+ }
+ } while ( NULL != ( p = rCode.NextRPN() ) );
+ if( isGroundWater() )
{
- ocl_calc.CreateBuffer32Bits(fpOclSrcData,npOclStartPos,npOclEndPos,rowSize);
- ocl_calc.CreateBuffer32Bits(fpLeftData,fpRightData,rowSize);
+ isSample = true;
}
- //printf("pptrr is %d,%d,%d\n",fpOclSrcData,npOclStartPos,npOclEndPos);
- }
-///////////////////////////////////////////////////////////////////////////////////////////
-
- // 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())
+ mnOperatorCount = 0;
+ mnSingleCount = 0;
+ mnSvDoubleCount = 0;
+ p = rCode.FirstRPN();
+ if(isSample)
{
- switch (p->GetType())
+ do
{
- case formula::svSingleVectorRef:
+ if ( ocPush == p->GetOpCode() && formula::svDouble == p->GetType() )
{
- 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);
+ mdpSvdouble[mnSvDoubleCount++] = p->GetDouble();
}
- break;
- case formula::svDoubleVectorRef:
+ else if( ocPush == p->GetOpCode() && formula::svDoubleVectorRef == p->GetType())
{
- 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;
- //store the a matix`s rowsize and colsize,use it to calculate the matix`s size
- ocl_calc.nFormulaRowSize = nRowSize;
- ocl_calc.nFormulaColSize = nColSize;
- ScMatrixRef pMat(new ScMatrix(nColSize, nRowSize, 0.0));
- if(ocl_calc.GetOpenclState())
- {
- npOclStartPos[i] = nRowStart; // record the start position
- npOclEndPos[i] = nRowEnd; // record the end position
- }
- int nTempOpcode;
- const formula::FormulaToken* pTemp = p;
- pTemp=aCode2.Next();
- nTempOpcode=pTemp->GetOpCode();
- while(1)
- {
- nTempOpcode=pTemp->GetOpCode();
- if(nTempOpcode!=ocOpen && nTempOpcode!=ocPush)
- break;
- pTemp=aCode2.Next();
- }
- if((!aIsAlloc) && (ocl_calc.GetOpenclState())&& (nTempOpcode == ocSumProduct))
- {
- //nColSize * rowSize is the data size , but except the the head of data will use less the nRowSize
- //the other all use nRowSize times . and it must aligen so add nRowSize-1.
- nSumproductSize = nRowSize+nColSize * rowSize*nRowSize-1;
- if(ocl_calc.gpuEnv.mnKhrFp64Flag==1 || ocl_calc.gpuEnv.mnAmdFp64Flag==1)
- ocl_calc.CreateBuffer64Bits(dpSumProMergeLfData,dpSumProMergeRtData,npSumSize,nSumproductSize,rowSize);
- else
- ocl_calc.CreateBuffer32Bits(fpSumProMergeLfData,fpSumProMergeRtData,npSumSize,nSumproductSize,rowSize);
- aIsAlloc = true;
- isSumProduct=true;
- }
- if(isSumProduct)
- {
- if(ocl_calc.gpuEnv.mnKhrFp64Flag==1 || ocl_calc.gpuEnv.mnAmdFp64Flag==1)
- {
- if(nCountMatix%2==0)
- dpSaveData = dpSumProMergeLfData;
- else
- dpSaveData = dpSumProMergeRtData;
- }
- else
- {
- if(nCountMatix%2==0)
- fpSaveData = fpSumProMergeLfData;
- else
- fpSaveData = fpSumProMergeRtData;
- }
- }
- for (size_t nCol = 0; nCol < nColSize; ++nCol)
+ const formula::DoubleVectorRefToken* pDvr = static_cast< const formula::DoubleVectorRefToken* >( p );
+ const std::vector< const double* >& rArrays = pDvr->GetArrays();
+ uint rArraysSize = rArrays.size();
+ int nMoreColSize = 0;
+ DoubleVectorFormula *SvDoubleTemp = new DoubleVectorFormula();
+ if( rArraysSize > 1 )
{
- const double* pArray = rArrays[nCol];
- if( NULL==pArray )
- {
- fprintf(stderr,"Error: pArray is NULL!\n");
- free(rResult);
- return false;
- }
- if(ocl_calc.GetOpenclState())
- {
- for( size_t u=nRowStart; u<=nRowEnd; u++ )
- {
- if(ocl_calc.gpuEnv.mnKhrFp64Flag==1 || ocl_calc.gpuEnv.mnAmdFp64Flag==1)
- {
- dpOclSrcData[u] = pArray[u];
- //fprintf(stderr,"dpOclSrcData[%d] is %f.\n",u,dpOclSrcData[u]);
- if(isSumProduct)
- dpSaveData[u+nRowSize*nCol + nRowStart* nColSize * nRowSize-nRowStart] = pArray[u];
- }
- else
- {
- // Many video cards can't support double type in kernel, so need transfer the double to float
- fpOclSrcData[u] = (float)pArray[u];
- //fprintf(stderr,"fpOclSrcData[%d] is %f.\n",u,fpOclSrcData[u]);
- if(isSumProduct)
- fpSaveData[u+nRowSize*nCol + nRowStart* nColSize * nRowSize-nRowStart] = (float)pArray[u];
- }
- }
- }
-
- for (size_t nRow = 0; nRow < nRowSize; ++nRow)
+ double *dpMoreColData = NULL;
+ for ( uint loop=0; loop < rArraysSize; loop++ )
{
- if (nRowStart + nRow < p2->GetArrayLength())
+ dpOclSrcData = rArrays[loop];
+ nSrcDataSize = pDvr->GetArrayLength();
+ nMoreColSize += nSrcDataSize;
+ dpMoreColData = (double *) realloc(dpMoreColData,nMoreColSize * sizeof(double));
+ for ( uint j = nMoreColSize - nSrcDataSize, i = 0; i < nSrcDataSize; i++, j++ )
{
- double fVal = pArray[nRowStart+nRow];
- pMat->PutDouble(fVal, nCol, nRow);
+ dpMoreColData[j] = dpOclSrcData[i];
}
}
+ dpOclSrcData = dpMoreColData;
+ nSrcDataSize = nMoreColSize;
}
-
- ScMatrixToken aTok(pMat);
- aCode2.AddToken(aTok);
- if(isSumProduct)
+ else
{
- npSumSize[nCountMatix/2] =nRowSize*nColSize;
- nCountMatix++;
+ dpOclSrcData = rArrays[0];
+ nSrcDataSize = pDvr->GetArrayLength();
+ SvDoubleTemp->mdpInputData = dpOclSrcData;
+ SvDoubleTemp->mnInputDataSize = nSrcDataSize;
+ SvDoubleTemp->mnInputStartPosition = mnpOclStartPos[nCountNum*mnRowSize];
+ SvDoubleTemp->mnInputEndPosition = mnpOclEndPos[nCountNum*mnRowSize];
+ SvDoubleTemp->mnInputStartOffset = mnpOclStartPos[nCountNum*mnRowSize+1]-mnpOclStartPos[nCountNum*mnRowSize];
+ SvDoubleTemp->mnInputEndOffset = mnpOclEndPos[nCountNum*mnRowSize+1]-mnpOclEndPos[nCountNum*mnRowSize];
+ mDoubleArray[mnDoubleCount++] = SvDoubleTemp;
+ nCountNum++;
}
}
- break;
- default:
- aCode2.AddToken(*p);
- }
+ else if( ocPush == p->GetOpCode() && formula::svSingleVectorRef == p->GetType() )
+ {
+ const formula::SingleVectorRefToken* pSvr = static_cast<const formula::SingleVectorRefToken*>( p );
+ dpBinaryData = pSvr->GetArray();
+ uint nArrayLen = pSvr->GetArrayLength();
+ SingleVectorFormula *SignleTemp = new SingleVectorFormula() ;
+ if(isSingle)
+ {
+ SignleTemp = mSingleArray[--mnSingleCount];
+ SignleTemp->mdpInputRightData = dpBinaryData;
+ SignleTemp->mnInputRightDataSize = nArrayLen;
+ SignleTemp->mnInputRightStartPosition = 0;
+ SignleTemp->mnInputRightOffset = 0;
+ isSingle = false;
+ }
+ else
+ {
+ SignleTemp = new SingleVectorFormula();
+ SignleTemp->mdpInputLeftData = dpBinaryData;
+ SignleTemp->mnInputLeftDataSize = nArrayLen;
+ SignleTemp->mdpInputRightData = NULL;
+ SignleTemp->mnInputRightDataSize = 0;
+ SignleTemp->mnInputLeftStartPosition = 0;
+ SignleTemp->mnInputLeftOffset = 0;
+ isSingle = true;
+ }
+ mSingleArray[mnSingleCount++] = SignleTemp;
+ }
+ else
+ {
+ nOclOp = p->GetOpCode();
+ mnOperatorGroup[mnOperatorCount++] = nOclOp;
+ }
+ } while ( NULL != ( p = rCode.NextRPN() ) );
+ if ( !chooseFunction( ocl_calc, pResult ) )
+ return false;
+ else
+ dpResult = pResult;
}
-
- ScFormulaCell* pDest = rDoc.GetFormulaCell(aTmpPos);
- if (!pDest)
- return false;
- if(ocl_calc.GetOpenclState())
+ else
{
- const formula::FormulaToken *pCur = aCode2.First();
- aCode2.Reset();
- while( ( pCur = aCode2.Next() ) != NULL )
+ agency aChooseAction;
+
+ do
{
- OpCode eOp = pCur->GetOpCode();
- if(eOp==0)
+ if ( ocPush == p->GetOpCode() && formula::svDouble == p->GetType() )
+ {
+ dpSvDouble = (double *) malloc( sizeof(double ) * mnRowSize );
+ double dTempValue = p->GetDouble();
+ for ( uint i = 0; i < mnRowSize; i++ )
+ dpSvDouble[i] = dTempValue;
+ srdDataPush( new SourceData( dpSvDouble, mnRowSize ) );
+ collectDoublePointers( dpSvDouble );
+ }
+ else if( ocPush == p->GetOpCode() && formula::svDoubleVectorRef == p->GetType())
{
- if(ocl_calc.gpuEnv.mnKhrFp64Flag==1 || ocl_calc.gpuEnv.mnAmdFp64Flag==1)
+ const formula::DoubleVectorRefToken* pDvr = static_cast< const formula::DoubleVectorRefToken* >( p );
+ const std::vector< const double* >& rArrays = pDvr->GetArrays();
+ unsigned int rArraysSize = rArrays.size();
+ int nMoreColSize = 0;
+ if(rArraysSize > 1)
{
- if(nCount3%2==0)
- dpLeftData[nCount1++] = pCur->GetDouble();
- else
- dpRightData[nCount2++] = pCur->GetDouble();
- nCount3++;
+ double *dpMoreColData = NULL;
+ for( uint loop=0; loop < rArraysSize; loop++ )
+ {
+ dpOclSrcData = rArrays[loop];
+ nSrcDataSize = pDvr->GetArrayLength();
+ nMoreColSize += nSrcDataSize;
+ dpMoreColData = (double *) realloc(dpMoreColData,nMoreColSize * sizeof(double));
+ for(uint j=nMoreColSize-nSrcDataSize,i=0;i<nSrcDataSize;i++,j++)
+ {
+ dpMoreColData[j] = dpOclSrcData[i];
+ }
+ }
+ dpOclSrcData = dpMoreColData;
+ nSrcDataSize = nMoreColSize;
+ collectDoublePointers( dpMoreColData );
}
else
{
- if(nCount3%2==0)
- fpLeftData[nCount1++] = (float)pCur->GetDouble();
- else
- fpRightData[nCount2++] = (float)pCur->GetDouble();
- nCount3++;
+ dpOclSrcData = rArrays[0];
+ nSrcDataSize = pDvr->GetArrayLength();
}
+ srdDataPush( new SourceData( dpOclSrcData,nSrcDataSize,rArraysSize ) );
}
- else if( eOp!=ocOpen && eOp!=ocClose &&eOp != ocSep)
- nOclOp = eOp;
-
-// if(count1>0){//dbg
-// fprintf(stderr,"leftData is %f.\n",fpLeftData[count1-1]);
-// count1--;
-// }
-// if(count2>0){//dbg
-// fprintf(stderr,"rightData is %f.\n",fpRightData[count2-1]);
-// count2--;
-// }
- }
- }
-
- if(!getenv("SC_GPU")||!ocl_calc.GetOpenclState())
- {
- //fprintf(stderr,"ccCPU flow...\n\n");
- generateRPNCode(rDoc, aTmpPos, aCode2);
- ScInterpreter aInterpreter(pDest, &rDoc, aTmpPos, aCode2);
- aInterpreter.Interpret();
- pDest->SetResultToken(aInterpreter.GetResultToken().get());
- pDest->ResetDirty();
- pDest->SetChanged(true);
+ else if( ocPush == p->GetOpCode() && formula::svSingleVectorRef == p->GetType() )
+ {
+ const formula::SingleVectorRefToken* pSvr = static_cast<const formula::SingleVectorRefToken*>( p );
+ dpBinaryData = pSvr->GetArray();
+ nSrcDataSize = pSvr->GetArrayLength();
+ srdDataPush( new SourceData( dpBinaryData, nSrcDataSize ) );
+ }
+ else
+ {
+ nOclOp = p->GetOpCode();
+ aChooseAction.calculate(nOclOp,mnRowSize,ocl_calc,mnpOclStartPos,mnpOclEndPos,this);
+ mnSingleCount = 0;
+ mnDoubleCount = 0;
+ mnSvDoubleCount = 0;
+ mnOperatorCount = 0;
+ mnPositonLen = 0;
+ }
+ } while ( NULL != ( p = rCode.NextRPN() ) );
+ SourceData * sResult = srdDataPop();
+ dpResult = sResult->getDouleData();
}
- } // for loop end (xGroup->mnLength)
-
- // For GPU calculation
- if(getenv("SC_GPU")&&ocl_calc.GetOpenclState())
- {
- fprintf(stderr,"ggGPU flow...\n\n");
- printf(" oclOp is... %d\n",nOclOp);
- osl_getSystemTime(&aTimeBefore); //timer
- if(ocl_calc.gpuEnv.mnKhrFp64Flag==1 || ocl_calc.gpuEnv.mnAmdFp64Flag==1)
+ rDoc.SetFormulaResults( rTopPos, dpResult, mnRowSize );
+ freeDoublePointers();
+ if ( pResult )
{
- fprintf(stderr,"ggGPU double precision flow...\n\n");
- //double precision
- switch(nOclOp)
- {
- case ocAdd:
- ocl_calc.OclHostArithmeticOperator64Bits("oclSignedAdd",dpLeftData,dpRightData,rResult,nCount1);
- break;
- case ocSub:
- ocl_calc.OclHostArithmeticOperator64Bits("oclSignedSub",dpLeftData,dpRightData,rResult,nCount1);
- break;
- case ocMul:
- ocl_calc.OclHostArithmeticOperator64Bits("oclSignedMul",dpLeftData,dpRightData,rResult,nCount1);
- break;
- case ocDiv:
- ocl_calc.OclHostArithmeticOperator64Bits("oclSignedDiv",dpLeftData,dpRightData,rResult,nCount1);
- break;
- case ocMax:
- ocl_calc.OclHostFormulaStatistics64Bits("oclFormulaMax",dpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocMin:
- ocl_calc.OclHostFormulaStatistics64Bits("oclFormulaMin",dpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocAverage:
- ocl_calc.OclHostFormulaStatistics64Bits("oclFormulaAverage",dpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocSum:
- ocl_calc.OclHostFormulaStatistics64Bits("oclFormulaSum",dpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocCount:
- ocl_calc.OclHostFormulaCount64Bits(npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocSumProduct:
- ocl_calc.OclHostFormulaSumProduct64Bits(dpSumProMergeLfData,dpSumProMergeRtData,npSumSize,rResult,rowSize);
- break;
- default:
- fprintf(stderr,"No OpenCL function for this calculation.\n");
- break;
- }
+ free( pResult );
+ pResult = NULL;
}
- else
+ if ( mnpOclStartPos )
{
- fprintf(stderr,"ggGPU float precision flow...\n\n");
- //float precision
- switch(nOclOp)
- {
- case ocAdd:
- ocl_calc.OclHostArithmeticOperator32Bits("oclSignedAdd",fpLeftData,fpRightData,rResult,nCount1);
- break;
- case ocSub:
- ocl_calc.OclHostArithmeticOperator32Bits("oclSignedSub",fpLeftData,fpRightData,rResult,nCount1);
- break;
- case ocMul:
- ocl_calc.OclHostArithmeticOperator32Bits("oclSignedMul",fpLeftData,fpRightData,rResult,nCount1);
- break;
- case ocDiv:
- ocl_calc.OclHostArithmeticOperator32Bits("oclSignedDiv",fpLeftData,fpRightData,rResult,nCount1);
- break;
- case ocMax:
- ocl_calc.OclHostFormulaStatistics32Bits("oclFormulaMax",fpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocMin:
- ocl_calc.OclHostFormulaStatistics32Bits("oclFormulaMin",fpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocAverage:
- ocl_calc.OclHostFormulaStatistics32Bits("oclFormulaAverage",fpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocSum:
- ocl_calc.OclHostFormulaStatistics32Bits("oclFormulaSum",fpOclSrcData,npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocCount:
- ocl_calc.OclHostFormulaCount32Bits(npOclStartPos,npOclEndPos,rResult,rowSize);
- break;
- case ocSumProduct:
- ocl_calc.OclHostFormulaSumProduct32Bits(fpSumProMergeLfData,fpSumProMergeRtData,npSumSize,rResult,rowSize);
- break;
- default:
- fprintf(stderr,"No OpenCL function for this calculation.\n");
- break;
- }
+ free( mnpOclStartPos );
+ mnpOclStartPos = NULL;
}
-
- /////////////////////////////////////////////////////
- osl_getSystemTime(&aTimeAfter);
- double diff = getTimeDiff(aTimeAfter, aTimeBefore);
- //if (diff >= 1.0)
+ if ( mnpOclEndPos )
{
- fprintf(stderr,"OpenCL,diff...%f.\n",diff);
+ free( mnpOclEndPos );
+ mnpOclEndPos = NULL;
}
-/////////////////////////////////////////////////////
-
-//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(rResult)
- free(rResult);
-
- return true;
+ return true;
+ } // getOpenclState() End
+ else
+ return false;
}
/// Special case of formula compiler for groundwatering
@@ -487,11 +964,11 @@ public:
FormulaGroupInterpreterSoftware()
{
fprintf(stderr,"\n\n ***** Groundwater Backend *****\n\n\n");
- OclCalc::InitEnv();
+ OclCalc::initEnv();
}
virtual ~FormulaGroupInterpreterGroundwater()
{
- OclCalc::ReleaseOpenclRunEnv();
+ OclCalc::releaseOpenclRunEnv();
}
virtual ScMatrixRef inverseMatrix(const ScMatrix& /* rMat */) { return ScMatrixRef(); }
@@ -567,7 +1044,7 @@ bool FormulaGroupInterpreterGroundwater::interpretCL(ScDocument& rDoc, const ScA
fprintf (stderr, "Calculate !");
- double *pResult = ocl_calc.OclSimpleDeltaOperation( eOp, pGroundWaterDataArray,
+ double *pResult = ocl_calc.oclSimpleDeltaOperation( eOp, pGroundWaterDataArray,
pArrayToSubtractOneElementFrom,
(size_t) xGroup->mnLength, delta );
RETURN_IF_FAIL(pResult != NULL, "buffer alloc / calculaton failed");
diff --git a/sc/source/core/opencl/oclkernels.hxx b/sc/source/core/opencl/oclkernels.hxx
index bcd7db0..53917b3 100644
--- a/sc/source/core/opencl/oclkernels.hxx
+++ b/sc/source/core/opencl/oclkernels.hxx
@@ -7,8 +7,8 @@
* file, You can obtain one at http://mozilla.org/MPL/2.0/.
*/
-#ifndef _OCL_KERNEL_H_
-#define _OCL_KERNEL_H_
+#ifndef SC_OCLKERNELS_HXX
+#define SC_OCLKERNELS_HXX
#ifndef USE_EXTERNAL_KERNEL
#define KERNEL( ... )# __VA_ARGS__
@@ -24,6 +24,97 @@ const char *kernel_src = KERNEL(
\n#pragma OPENCL EXTENSION cl_amd_fp64 : enable\n
\n#else\n
\n#endif\n
+inline fp_t oclAverage( const uint id,__global fp_t *values,__global uint *startArray,__global uint *endArray)
+{
+ uint start = startArray[id];
+ uint end = endArray[id];
+ fp_t fSum = 0.0;
+ fp_t zero[16] = {0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f,0.0f};
+ fp_t16 vSum=vload16(0,zero);
+ fp_t16 ptr;
+ __global fp_t *p = values;
+ p+= start;
+
+ for(int i = 0; i < (end - start + 1)/16; ++i)
+ {
+ ptr=vload16(0,p);
+ vSum += ptr;
+ p+=16;
+ }
+ int lastData = (end-start+1)%16;
+ for(int i = 0; i <lastData; i++)
+ {
+ fSum += *p;
+ p+=1;
+ }
+ vSum.s01234567 = vSum.s01234567+vSum.s89abcdef;
+ vSum.s0123 = vSum.s0123+vSum.s4567;
+ vSum.s01 = vSum.s01+vSum.s23;
+ vSum.s0 = vSum.s0+vSum.s1;
+ fSum = vSum.s0+fSum;
+ fp_t fVal = fSum/(end-start+1);
+ return fVal;
+}
+inline fp_t oclMax( const uint id,__global fp_t *values,__global uint *startArray,__global uint *endArray)
+{
+ uint start = startArray[id];
+ uint end = endArray[id];
+ fp_t fMax = values[start];
+ fp_t zero[16] = {fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax,fMax};
+ fp_t16 vMax=vload16(0,zero);
+ //Max
+ fp_t16 ptr;
+ __global fp_t *p = values;
+ p+= start;
+ for(int i = 0; i < (end - start + 1)/16; ++i)
+ {
+ ptr=vload16(0,p);
+ vMax = fmax(vMax,ptr);
+ p+=16;
+ }
+ int lastData = (end-start+1)%16;
+ for(int i = 0; i <lastData; i++)
+ {
+ fMax = fmax(fMax,*p);
+ p+=1;
+ }
+ vMax.s01234567 = fmax(vMax.s01234567, vMax.s89abcdef);
+ vMax.s0123 = fmax(vMax.s0123, vMax.s4567);
+ vMax.s01 = fmax(vMax.s01, vMax.s23);
+ vMax.s0 = fmax(vMax.s0, vMax.s1);
+ fMax = fmax(vMax.s0, fMax);
+ return fMax;
+}
+inline fp_t oclMin( const uint id,__global fp_t *values,__global uint *startArray,__global uint *endArray)
+{
+ uint start = startArray[id];
+ uint end = endArray[id];
+ fp_t fMin = values[start];
+ fp_t zero[16] = {fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin,fMin};
+ fp_t16 vMin=vload16(0,zero);
+ //Min
+ fp_t16 ptr;
+ __global fp_t *p = values;
+ p+= start;
+ for(int i = 0; i < (end - start + 1)/16; ++i)
+ {
+ ptr=vload16(0,p);
+ vMin = fmin(vMin,ptr);
+ p+=16;
+ }
+ int lastData = (end-start+1)%16;
+ for(int i = 0; i <lastData; i++)
+ {
+ fMin = fmin(fMin,*p);
+ p+=1;
+ }
+ vMin.s01234567 = fmin(vMin.s01234567, vMin.s89abcdef);
+ vMin.s0123 = fmin(vMin.s0123, vMin.s4567);
+ vMin.s01 = fmin(vMin.s01, vMin.s23);
+ vMin.s0 = fmin(vMin.s0, vMin.s1);
+ fMin = fmin(vMin.s0, fMin);
+ return fMin;
+}
__kernel void oclSignedAdd(__global fp_t *ltData,__global fp_t *rtData,__global fp_t *otData)
{
@@ -31,7 +122,6 @@ __kernel void oclSignedAdd(__global fp_t *ltData,__global fp_t *rtData,__global
otData[id] = ltData[id] + rtData[id];
}
-
__kernel void oclSignedSub(__global fp_t *ltData,__global fp_t *rtData,__global fp_t *otData)
{
const unsigned int id = get_global_id(0);
@@ -41,39 +131,31 @@ __kernel void oclSignedSub(__global fp_t *ltData,__global fp_t *rtData,__global
__kernel void oclSignedMul(__global fp_t *ltData,__global fp_t *rtData,__global fp_t *otData)
{
int id = get_global_id(0);
- otData[id] =ltData[id] * rtData[id];
+ otData[id] = ltData[id] * rtData[id];
}
__kernel void oclSignedDiv(__global fp_t *ltData,__global fp_t *rtData,__global fp_t *otData)
{
const unsigned int id = get_global_id(0);
- otData[id] = ltData[id] / rtData[id];
+ fp_t divisor = rtData[id];
+ if ( divisor != 0 )
+ otData[id] = ltData[id] / divisor;
+ else
+ otData[id] = 0.0;
}
__kernel void oclFormulaMin(__global fp_t *input,__global uint *start,__global uint *end,__global fp_t *output)
{
const unsigned int id = get_global_id(0);
- unsigned int startFlag = start[id];
- unsigned int endFlag = end[id];
- fp_t fMinVal = input[startFlag];
- for(int i=startFlag;i<=endFlag;i++)
- {
- fMinVal = fmin( fMinVal, input[i] );
- }
- output[id] = fMinVal;
+ fp_t fVal = oclMin(id,input,start,end);
+ output[id] = fVal ;
}
__kernel void oclFormulaMax(__global fp_t *input,__global uint *start,__global uint *end,__global fp_t *output)
{
const unsigned int id = get_global_id(0);
- unsigned int startFlag = start[id];
- unsigned int endFlag = end[id];
- fp_t fMaxVal = input[startFlag];
- for ( int i = startFlag; i <= endFlag; i++ )
- {
- fMaxVal = fmax( fMaxVal, input[i] );
- }
- output[id] = fMaxVal;
+ fp_t fVal = oclMax(id,input,start,end);
+ output[id] = fVal ;
}
//Sum
__kernel void oclFormulaSum(__global fp_t *input,__global uint *start,__global uint *end,__global fp_t *output)
@@ -94,12 +176,10 @@ __kernel void oclFormulaCount(__global uint *start,__global uint *end,__global f
__kernel void oclFormulaAverage(__global fp_t *input,__global uint *start,__global uint *end,__global fp_t *output)
{
const unsigned int id = get_global_id(0);
- fp_t sum=0.0;
- for(int i = start[id];i<=end[id];i++)
- sum += input[i];
- output[id] = sum / (end[id]-start[id]+1);
-}
+ fp_t fVal = oclAverage(id,input,start,end);
+ output[id] = fVal ;
+}
//Sumproduct
__kernel void oclFormulaSumproduct(__global fp_t *firstCol,__global uint* npSumSize,__global fp_t *output,uint nMatixSize)
{
@@ -147,7 +227,7 @@ __kernel void oclMinDelta(__global fp_t *values, __global fp_t *subtract, uint s
// Min
fp_t fMinVal = values[start];
- for(int i=start+1;i < end;i++)
+ for ( int i = start + 1; i < end; i++ )
{
if(values[i]<fMinVal)
fMinVal = values[i];
@@ -177,14 +257,14 @@ __kernel void oclFormulaMtxInv(__global fp_t * fpMatrixInput, __global fp_t * fp
fpP[nOffset*nDimension+nId] = fpP[nMax*nDimension+nId];
fpP[nMax*nDimension+nId] = dMovebuffer;
}
-__kernel void oclMatrixSolve(__global fp_t * fpMatrixInput,__global fp_t * fpMatrixOutput,__global fp_t * fpP,__global fp_t * fpY)
+__kernel void oclMatrixSolve(__global fp_t * fpMatrixInput,__global fp_t * fpMatrixOutput,__global fp_t * fpP,__global fp_t * fpY,__global uint* npDim)
{
int nId = get_global_id(0);
- int nDimension = get_global_size(0);
-
+ int nDimension = npDim[nId];
+ fp_t fsum = 0.0;
for ( int yi=0; yi < nDimension; yi++ )
{
- fp_t fsum = 0.0;
+ fsum = 0.0;
for ( int yj=0; yj < nDimension; yj++ )
{
fsum += fpMatrixInput[yi*nDimension+yj] * fpY[nId+yj*nDimension];
@@ -194,7 +274,7 @@ __kernel void oclMatrixSolve(__global fp_t * fpMatrixInput,__global fp_t * fpMat
}
for ( int xi = nDimension - 1; xi >= 0; xi-- )
{
- fp_t fsum = 0.0;
+ fsum = 0.0;
for ( int xj = 0; xj < nDimension; xj++ )
{
fsum += fpMatrixInput[xi*nDimension+xj] * fpMatrixOutput[nId+nDimension*xj];
@@ -203,6 +283,101 @@ __kernel void oclMatrixSolve(__global fp_t * fpMatrixInput,__global fp_t * fpMat
}
}
+__kernel void oclAverageAdd(__global fp_t *values,__global fp_t *addend, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fVal = oclAverage(id,values,startArray,endArray);
+ output[id] = fVal + addend[id];
+}
+
+__kernel void oclAverageSub(__global fp_t *values,__global fp_t *subtract, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fVal = oclAverage(id,values,startArray,endArray);
+ output[id] = fVal - subtract[id];
+}
+
+__kernel void oclAverageMul(__global fp_t *values,__global fp_t *multiplier, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fVal = oclAverage(id,values,startArray,endArray);
+ output[id] = fVal * multiplier[id];
+}
+__kernel void oclAverageDiv(__global fp_t *values,__global fp_t *div, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fVal = oclAverage(id,values,startArray,endArray);
+ fp_t divisor = div[id];
+ if ( divisor != 0 )
+ output[id] = fVal / divisor;
+ else
+ output[id] = 0.0;
+}
+
+__kernel void oclMinAdd(__global fp_t *values, __global fp_t *addend, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMin = oclMin(id,values,startArray,endArray);
+ output[id] = fMin + addend[id];
+}
+
+__kernel void oclMinSub(__global fp_t *values, __global fp_t *subtract, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMin = oclMin(id,values,startArray,endArray);
+ output[id] = fMin - subtract[id];
+}
+__kernel void oclMinMul(__global fp_t *values, __global fp_t *multiplier, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMin = oclMin(id,values,startArray,endArray);
+ output[id] = fMin * multiplier[id];
+}
+__kernel void oclMinDiv(__global fp_t *values, __global fp_t *div, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMin = oclMin(id,values,startArray,endArray);
+ fp_t divisor = div[id];
+ if ( divisor != 0 )
+ output[id] = fMin / divisor;
+ else
+ output[id] = 0.0;
+}
+__kernel void oclMaxAdd(__global fp_t *values, __global fp_t *addend, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMax = oclMax(id,values,startArray,endArray);
+ output[id] = fMax + addend[id];
+}
+
+__kernel void oclMaxSub(__global fp_t *values, __global fp_t *subtract, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMax = oclMax(id,values,startArray,endArray);
+ output[id] = fMax - subtract[id];
+}
+__kernel void oclMaxMul(__global fp_t *values, __global fp_t *multiplier, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMax = oclMax(id,values,startArray,endArray);
+ output[id] = fMax * multiplier[id];
+}
+__kernel void oclMaxDiv(__global fp_t *values, __global fp_t *div, __global uint *startArray, __global uint *endArray, __global fp_t *output)
+{
+ const unsigned int id = get_global_id(0);
+ fp_t fMax = oclMax(id,values,startArray,endArray);
+ fp_t divisor = div[id];
+ if ( divisor != 0 )
+ output[id] = fMax / divisor;
+ else
+ output[id] = 0.0;
+}
+
+__kernel void oclSub( fp_t ltData, __global fp_t *rtData, __global fp_t *outData )
+{
+ const unsigned int id = get_global_id(0);
+ outData[id] = ltData - rtData[id];
+}
);
#endif // USE_EXTERNAL_KERNEL
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 0557717..6db498b 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -7,19 +7,19 @@
* file, You can obtain one at http://mozilla.org/MPL/2.0/.
*/
-#include "openclwrapper.hxx"
-
+#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <cmath>
#include "sal/config.h"
#include "random.hxx"
+#include "openclwrapper.hxx"
#include "oclkernels.hxx"
#ifdef SAL_WIN32
#include <Windows.h>
#endif
-//#define USE_KERNEL_FILE
-
+//#define USE_MAP_BUFFER
+using namespace std;
GPUEnv OpenclDevice::gpuEnv;
int OpenclDevice::isInited =0;
@@ -44,9 +44,9 @@ int OpenclDevice::isInited =0;
HINSTANCE HOpenclDll = NULL;
void * OpenclDll = NULL;
-int OpenclDevice::LoadOpencl()
+int OpenclDevice::loadOpencl()
{
- //fprintf(stderr, " LoadOpenclDllxx... \n");
+ //fprintf(stderr, " loadOpenclDllxx... \n");
OpenclDll = static_cast<HINSTANCE>( HOpenclDll );
OpenclDll = LoadLibrary( OPENCL_DLL_NAME );
if ( !static_cast<HINSTANCE>( OpenclDll ) )
@@ -59,7 +59,7 @@ int OpenclDevice::LoadOpencl()
return OCLSUCCESS;
}
-void OpenclDevice::FreeOpenclDll()
+void OpenclDevice::freeOpenclDll()
{
fprintf(stderr, " Free opencllo.dll ... \n");
if ( !static_cast<HINSTANCE>( OpenclDll ) )
@@ -67,39 +67,39 @@ void OpenclDevice::FreeOpenclDll()
}
#endif
-int OpenclDevice::InitEnv()
+int OpenclDevice::initEnv()
{
#ifdef SAL_WIN32
while( 1 )
{
- if( 1 == LoadOpencl() )
+ if( 1 == loadOpencl() )
break;
}
#endif
- InitOpenclRunEnv( 0 );
+ initOpenclRunEnv( 0 );
return 1;
}
-int OpenclDevice::ReleaseOpenclRunEnv()
+int OpenclDevice::releaseOpenclRunEnv()
{
- ReleaseOpenclEnv( &gpuEnv );
+ releaseOpenclEnv( &gpuEnv );
#ifdef SAL_WIN32
- FreeOpenclDll();
+ freeOpenclDll();
#endif
return 1;
}
///////////////////////////////////////////////////////
///////////////////////////////////////////////////////
-inline int OpenclDevice::AddKernelConfig( int kCount, const char *kName )
+inline int OpenclDevice::addKernelConfig( int kCount, const char *kName )
{
if ( kCount < 1 )
- fprintf(stderr,"Error: ( KCount < 1 )" SAL_DETAIL_WHERE "AddKernelConfig\n" );
+ fprintf(stderr,"Error: ( KCount < 1 )" SAL_DETAIL_WHERE "addKernelConfig\n" );
strcpy( gpuEnv.mArrykernelNames[kCount-1], kName );
gpuEnv.mnKernelCount++;
return 0;
}
-int OpenclDevice::RegistOpenclKernel()
+int OpenclDevice::registOpenclKernel()
{
if ( !gpuEnv.mnIsUserCreated )
memset( &gpuEnv, 0, sizeof(gpuEnv) );
@@ -107,37 +107,58 @@ int OpenclDevice::RegistOpenclKernel()
gpuEnv.mnFileCount = 0; //argc;
gpuEnv.mnKernelCount = 0UL;
- AddKernelConfig( 1, (const char*) "oclFormulaMin" );
- AddKernelConfig( 2, (const char*) "oclFormulaMax" );
- AddKernelConfig( 3, (const char*) "oclFormulaSum" );
- AddKernelConfig( 4, (const char*) "oclFormulaCount" );
- AddKernelConfig( 5, (const char*) "oclFormulaAverage" );
- AddKernelConfig( 6, (const char*) "oclFormulaSumproduct" );
- AddKernelConfig( 7, (const char*) "oclFormulaMtxInv" );
-
- AddKernelConfig( 8, (const char*) "oclSignedAdd" );
- AddKernelConfig( 9, (const char*) "oclSignedSub" );
- AddKernelConfig( 10, (const char*) "oclSignedMul" );
- AddKernelConfig( 11, (const char*) "oclSignedDiv" );
- AddKernelConfig( 12, (const char*) "oclAverageDelta" );
- AddKernelConfig( 13, (const char*) "oclMaxDelta" );
- AddKernelConfig( 14, (const char*) "oclMinDelta" );
- AddKernelConfig( 15, (const char*) "oclSubDelta" );
- AddKernelConfig( 16, (const char*) "oclLUDecomposition" );
+ addKernelConfig( 1, (const char*) "oclFormulaMin" );
+ addKernelConfig( 2, (const char*) "oclFormulaMax" );
+ addKernelConfig( 3, (const char*) "oclFormulaSum" );
+ addKernelConfig( 4, (const char*) "oclFormulaCount" );
+ addKernelConfig( 5, (const char*) "oclFormulaAverage" );
+ addKernelConfig( 6, (const char*) "oclFormulaSumproduct" );
+ addKernelConfig( 7, (const char*) "oclFormulaMtxInv" );
+
+ addKernelConfig( 8, (const char*) "oclSignedAdd" );
+ addKernelConfig( 9, (const char*) "oclSignedSub" );
+ addKernelConfig( 10, (const char*) "oclSignedMul" );
+ addKernelConfig( 11, (const char*) "oclSignedDiv" );
+ addKernelConfig( 12, (const char*) "oclAverageDelta" );
+ addKernelConfig( 13, (const char*) "oclMaxDelta" );
+ addKernelConfig( 14, (const char*) "oclMinDelta" );
+ addKernelConfig( 15, (const char*) "oclSubDelta" );
+ addKernelConfig( 16, (const char*) "oclLUDecomposition" );
+ addKernelConfig( 17, (const char*) "oclAverageDeltaRPN" );
+ addKernelConfig( 18, (const char*) "oclMaxDeltaRPN" );
+ addKernelConfig( 19, (const char*) "oclMinDeltaRPN" );
+ addKernelConfig( 20, (const char*) "oclMoreColArithmeticOperator" );
+ addKernelConfig( 21, (const char*) "oclColumnH" );
+ addKernelConfig( 22, (const char*) "oclColumnL" );
+ addKernelConfig( 23, (const char*) "oclColumnN" );
+ addKernelConfig( 24, (const char*) "oclColumnJ" );
+ addKernelConfig( 25, (const char*) "oclMaxSub" );
+ addKernelConfig( 26, (const char*) "oclAverageSub" );
+ addKernelConfig( 27, (const char*) "oclMinSub" );
+ addKernelConfig( 28, (const char*) "oclMaxAdd" );
+ addKernelConfig( 29, (const char*) "oclAverageAdd" );
+ addKernelConfig( 30, (const char*) "oclMinAdd" );
+ addKernelConfig( 31, (const char*) "oclMaxMul" );
+ addKernelConfig( 32, (const char*) "oclAverageMul" );
+ addKernelConfig( 33, (const char*) "oclMinMul" );
+ addKernelConfig( 34, (const char*) "oclMaxDiv" );
+ addKernelConfig( 35, (const char*) "oclAverageDiv" );
+ addKernelConfig( 36, (const char*) "oclMinDiv" );
+ addKernelConfig( 37, (const char*) "oclSub" );// for svDouble type
return 0;
}
OpenclDevice::OpenclDevice()
{
- //InitEnv();
+ //initEnv();
}
OpenclDevice::~OpenclDevice()
{
- //ReleaseOpenclRunEnv();
+ //releaseOpenclRunEnv();
}
-int OpenclDevice::SetKernelEnv( KernelEnv *envInfo )
+int OpenclDevice::setKernelEnv( KernelEnv *envInfo )
{
envInfo->mpkContext = gpuEnv.mpContext;
envInfo->mpkCmdQueue = gpuEnv.mpCmdQueue;
@@ -146,12 +167,12 @@ int OpenclDevice::SetKernelEnv( KernelEnv *envInfo )
return 1;
}
-int OpenclDevice::CheckKernelName( KernelEnv *envInfo, const char *kernelName )
+int OpenclDevice::checkKernelName( KernelEnv *envInfo, const char *kernelName )
{
- //printf("CheckKernelName,total count of kernels...%d\n", gpuEnv.kernelCount);
+ //printf("checkKernelName,total count of kernels...%d\n", gpuEnv.kernelCount);
int kCount;
int nFlag = 0;
- for ( kCount = 0; kCount < gpuEnv.mnKernelCount; kCount++ )
+ for ( kCount=0; kCount < gpuEnv.mnKernelCount; kCount++ )
{
if ( strcasecmp( kernelName, gpuEnv.mArrykernelNames[kCount]) == 0 )
{
@@ -174,7 +195,7 @@ int OpenclDevice::CheckKernelName( KernelEnv *envInfo, const char *kernelName )
return 1;
}
-int OpenclDevice::ConvertToString( const char *filename, char **source )
+int OpenclDevice::convertToString( const char *filename, char **source )
{
int file_size;
size_t result;
@@ -211,7 +232,7 @@ int OpenclDevice::ConvertToString( const char *filename, char **source )
return 0;
}
-int OpenclDevice::BinaryGenerated( const char * clFileName, FILE ** fhandle )
+int OpenclDevice::binaryGenerated( const char * clFileName, FILE ** fhandle )
{
unsigned int i = 0;
cl_int clStatus;
@@ -219,11 +240,23 @@ int OpenclDevice::BinaryGenerated( const char * clFileName, FILE ** fhandle )
char *str = NULL;
FILE *fd = NULL;
cl_uint numDevices=0;
- clStatus = clGetDeviceIDs(gpuEnv.mpPlatformID, // platform
- CL_DEVICE_TYPE_GPU, // device_type
- 0, // num_entries
- NULL, // devices ID
- &numDevices);
+ if ( getenv("SC_OPENCLCPU") )
+ {
+ clStatus = clGetDeviceIDs(gpuEnv.mpPlatformID, // platform
+ CL_DEVICE_TYPE_CPU, // device_type for CPU device
+ 0, // num_entries
+ NULL, // devices ID
+ &numDevices);
+ }
+ else
+ {
+ clStatus = clGetDeviceIDs(gpuEnv.mpPlatformID, // platform
+ CL_DEVICE_TYPE_GPU, // device_type for GPU device
+ 0, // num_entries
+ NULL, // devices ID
+ &numDevices);
+ }
+ CHECK_OPENCL( clStatus, "clGetDeviceIDs" );
for ( i = 0; i < numDevices; i++ )
{
char fileName[256] = { 0 }, cl_name[128] = { 0 };
@@ -248,7 +281,7 @@ int OpenclDevice::BinaryGenerated( const char * clFileName, FILE ** fhandle )
}
-int OpenclDevice::WriteBinaryToFile( const char* fileName, const char* birary, size_t numBytes )
+int OpenclDevice::writeBinaryToFile( const char* fileName, const char* birary, size_t numBytes )
{
FILE *output = NULL;
output = fopen( fileName, "wb" );
@@ -264,7 +297,7 @@ int OpenclDevice::WriteBinaryToFile( const char* fileName, const char* birary, s
}
-int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * clFileName )
+int OpenclDevice::generatBinFromKernelSource( cl_program program, const char * clFileName )
{
unsigned int i = 0;
cl_int clStatus;
@@ -307,12 +340,6 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c
binaries[i] = (char*) malloc( sizeof(char) * binarySizes[i] );
if ( binaries[i] == NULL )
{
- for ( unsigned int j = 0; j < i ; j++)
- {
- if (binaries[j])
- free(binaries[j]);
- }
- free(binaries);
return 0;
}
}
@@ -343,7 +370,7 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c
cl_name[str - clFileName] = '\0';
sprintf( fileName, "./%s-%s.bin", cl_name, deviceName );
- if ( !WriteBinaryToFile( fileName, binaries[i], binarySizes[i] ) )
+ if ( !writeBinaryToFile( fileName, binaries[i], binarySizes[i] ) )
{
printf("opencl-wrapper: write binary[%s] failds\n", fileName);
return 0;
@@ -382,7 +409,7 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c
return 1;
}
-int OpenclDevice::InitOpenclAttr( OpenCLEnv * env )
+int OpenclDevice::initOpenclAttr( OpenCLEnv * env )
{
if ( gpuEnv.mnIsUserCreated )
return 1;
@@ -397,7 +424,7 @@ int OpenclDevice::InitOpenclAttr( OpenCLEnv * env )
return 0;
}
-int OpenclDevice::CreateKernel( char * kernelname, KernelEnv * env )
+int OpenclDevice::createKernel( char * kernelname, KernelEnv * env )
{
int clStatus;
@@ -407,13 +434,13 @@ int OpenclDevice::CreateKernel( char * kernelname, KernelEnv * env )
return clStatus != CL_SUCCESS ? 1 : 0;
}
-int OpenclDevice::ReleaseKernel( KernelEnv * env )
+int OpenclDevice::releaseKernel( KernelEnv * env )
{
int clStatus = clReleaseKernel( env->mpkKernel );
return clStatus != CL_SUCCESS ? 1 : 0;
}
-int OpenclDevice::ReleaseOpenclEnv( GPUEnv *gpuInfo )
+int OpenclDevice::releaseOpenclEnv( GPUEnv *gpuInfo )
{
int i = 0;
int clStatus = 0;
@@ -448,18 +475,18 @@ int OpenclDevice::ReleaseOpenclEnv( GPUEnv *gpuInfo )
return 1;
}
-int OpenclDevice::RunKernelWrapper( cl_kernel_function function, const char * kernelName, void **usrdata )
+int OpenclDevice::runKernelWrapper( cl_kernel_function function, const char * kernelName, void **usrdata )
{
- printf("oclwrapper:RunKernel_wrapper...\n");
- if ( RegisterKernelWrapper( kernelName, function ) != 1 )
+ printf("oclwrapper:runKernel_wrapper...\n");
+ if ( registerKernelWrapper( kernelName, function ) != 1 )
{
- fprintf(stderr, "Error:RunKernel_wrapper:RegisterKernelWrapper fail!\n");
+ fprintf(stderr, "Error:runKernel_wrapper:registerKernelWrapper fail!\n");
return -1;
}
- return ( RunKernel( kernelName, usrdata ) );
+ return ( runKernel( kernelName, usrdata ) );
}
-int OpenclDevice::CachedOfKernerPrg( const GPUEnv *gpuEnvCached, const char * clFileName )
+int OpenclDevice::cachedOfKernerPrg( const GPUEnv *gpuEnvCached, const char * clFileName )
{
int i;
for ( i = 0; i < gpuEnvCached->mnFileCount; i++ )
@@ -476,7 +503,7 @@ int OpenclDevice::CachedOfKernerPrg( const GPUEnv *gpuEnvCached, const char * cl
return 0;
}
-int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
+int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
{
cl_int clStatus = 0;
size_t length;
@@ -484,12 +511,12 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
const char *source;
size_t source_size[1];
int b_error, binary_status, binaryExisted, idx;
- cl_uint numDevices;
+ size_t numDevices;
cl_device_id *mpArryDevsID;
FILE *fd, *fd1;
const char* filename = "kernel.cl";
- fprintf(stderr, "CompileKernelFile ... \n");
- if ( CachedOfKernerPrg(gpuInfo, filename) == 1 )
+ fprintf(stderr, "compileKernelFile ... \n");
+ if ( cachedOfKernerPrg(gpuInfo, filename) == 1 )
{
return 1;
}
@@ -500,7 +527,7 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
source_size[0] = strlen( source );
binaryExisted = 0;
- if ( ( binaryExisted = BinaryGenerated( filename, &fd ) ) == 1 )
+ if ( ( binaryExisted = binaryGenerated( filename, &fd ) ) == 1 )
{
clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES,
sizeof(numDevices), &numDevices, NULL );
@@ -519,14 +546,12 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
b_error |= fseek( fd, 0, SEEK_SET ) < 0;
if ( b_error )
{
- free(mpArryDevsID);
return 0;
}
binary = (char*) malloc( length + 2 );
if ( !binary )
{
- free(mpArryDevsID);
return 0;
}
@@ -539,12 +564,7 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
// grab the handles to all of the devices in the context.
clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_DEVICES,
sizeof( cl_device_id ) * numDevices, mpArryDevsID, NULL );
- if (clStatus != CL_SUCCESS)
- {
- fprintf (stderr, "OpenCL error code is %d at " SAL_DETAIL_WHERE " when clGetContextInfo .\n", clStatus);
- free(binary);
- return 0;
- }
+ CHECK_OPENCL( clStatus, "clGetContextInfo" );
fprintf(stderr, "Create kernel from binary\n");
gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary( gpuInfo->mpContext,numDevices,
@@ -637,17 +657,17 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
strcpy( gpuEnv.mArryKnelSrcFile[idx], filename );
if ( binaryExisted == 0 )
- GeneratBinFromKernelSource( gpuEnv.mpArryPrograms[idx], filename );
+ generatBinFromKernelSource( gpuEnv.mpArryPrograms[idx], filename );
gpuInfo->mnFileCount += 1;
return 1;
}
-int OpenclDevice::GetKernelEnvAndFunc( const char *kernelName, KernelEnv *env, cl_kernel_function *function)
+int OpenclDevice::getKernelEnvAndFunc( const char *kernelName, KernelEnv *env, cl_kernel_function *function)
{
int i;
- //printf("----------------OpenclDevice::GetKernelEnvAndFunc\n");
+ //printf("----------------OpenclDevice::getKernelEnvAndFunc\n");
for ( i = 0; i < gpuEnv.mnKernelCount; i++ )
{
if ( strcasecmp( kernelName, gpuEnv.mArrykernelNames[i]) == 0 )
@@ -663,14 +683,14 @@ int OpenclDevice::GetKernelEnvAndFunc( const char *kernelName, KernelEnv *env, c
return 0;
}
-int OpenclDevice::RunKernel( const char *kernelName, void **userdata)
+int OpenclDevice::runKernel( const char *kernelName, void **userdata)
{
KernelEnv kEnv;
cl_kernel_function function;
int status;
memset( &kEnv, 0, sizeof( KernelEnv ) );
- status = GetKernelEnvAndFunc( kernelName, &kEnv, &function );
+ status = getKernelEnvAndFunc( kernelName, &kEnv, &function );
strcpy( kEnv.mckKernelName, kernelName );
if ( status == 1 )
{
@@ -681,7 +701,7 @@ int OpenclDevice::RunKernel( const char *kernelName, void **userdata)
return 0;
}
-int OpenclDevice::InitOpenclRunEnv( int argc )
+int OpenclDevice::initOpenclRunEnv( int argc )
{
int status = 0;
if ( MAX_CLKERNEL_NUM <= 0 )
@@ -693,9 +713,9 @@ int OpenclDevice::InitOpenclRunEnv( int argc )
if ( !isInited )
{
- RegistOpenclKernel();
+ registOpenclKernel();
//initialize devices, context, comand_queue
- status = InitOpenclRunEnv( &gpuEnv );
+ status = initOpenclRunEnv( &gpuEnv );
if ( status )
{
printf("init_opencl_env failed.\n");
@@ -711,30 +731,30 @@ int OpenclDevice::InitOpenclRunEnv( int argc )
if( gpuEnv.mnKhrFp64Flag )
{
printf("----use khr double type in kernel----\n");
- status = CompileKernelFile( &gpuEnv, "-D KHR_DP_EXTENSION -Dfp_t=double" );
+ status = compileKernelFile( &gpuEnv, "-D KHR_DP_EXTENSION -Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16" );
}
else if( gpuEnv.mnAmdFp64Flag )
{
printf("----use amd double type in kernel----\n");
- status = CompileKernelFile( &gpuEnv, "-D AMD_DP_EXTENSION -Dfp_t=double" );
+ status = compileKernelFile( &gpuEnv, "-D AMD_DP_EXTENSION -Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16" );
}
else
{
printf("----use float type in kernel----\n");
- status = CompileKernelFile( &gpuEnv, "-Dfp_t=float" );
+ status = compileKernelFile( &gpuEnv, "-Dfp_t=float -Dfp_t4=float4 -Dfp_t16=float16" );
}
if ( status == 0 || gpuEnv.mnKernelCount == 0 )
{
- printf("CompileKernelFile failed.\n");
+ printf("compileKernelFile failed.\n");
return 1;
}
- printf("CompileKernelFile successed.\n");
+ printf("compileKernelFile successed.\n");
isInited = 1;
}
return 0;
}
-int OpenclDevice::InitOpenclRunEnv( GPUEnv *gpuInfo )
+int OpenclDevice::initOpenclRunEnv( GPUEnv *gpuInfo )
{
size_t length;
cl_int clStatus;
@@ -784,13 +804,22 @@ int OpenclDevice::InitOpenclRunEnv( GPUEnv *gpuInfo )
//if( !strcmp( platformName, "Advanced Micro Devices, Inc." ))
{
gpuInfo->mpPlatformID = platforms[i];
-
- clStatus = clGetDeviceIDs(gpuInfo->mpPlatformID, // platform
- CL_DEVICE_TYPE_GPU, // device_type
- 0, // num_entries
- NULL, // devices
- &numDevices);
-
+ if ( getenv("SC_OPENCLCPU") )
+ {
+ clStatus = clGetDeviceIDs(gpuInfo->mpPlatformID, // platform
+ CL_DEVICE_TYPE_CPU, // device_type for CPU device
+ 0, // num_entries
+ NULL, // devices
+ &numDevices);
+ }
+ else
+ {
+ clStatus = clGetDeviceIDs(gpuInfo->mpPlatformID, // platform
+ CL_DEVICE_TYPE_GPU, // device_type for GPU device
+ 0, // num_entries
+ NULL, // devices
+ &numDevices);
+ }
if ( clStatus != CL_SUCCESS )
continue;
@@ -809,8 +838,15 @@ int OpenclDevice::InitOpenclRunEnv( GPUEnv *gpuInfo )
cps[0] = CL_CONTEXT_PLATFORM;
cps[1] = (cl_context_properties) gpuInfo->mpPlatformID;
cps[2] = 0;
- // Check for GPU.
- gpuInfo->mDevType = CL_DEVICE_TYPE_GPU;
+ // Set device type for OpenCL
+ if ( getenv("SC_OPENCLCPU") )
+ {
+ gpuInfo->mDevType = CL_DEVICE_TYPE_CPU;
+ }
+ else
+ {
+ gpuInfo->mDevType = CL_DEVICE_TYPE_GPU;
+ }
gpuInfo->mpContext = clCreateContextFromType( cps, gpuInfo->mDevType, NULL, NULL, &clStatus );
if ( ( gpuInfo->mpContext == (cl_context) NULL) || ( clStatus != CL_SUCCESS ) )
@@ -877,10 +913,10 @@ int OpenclDevice::InitOpenclRunEnv( GPUEnv *gpuInfo )
return 0;
}
-int OpenclDevice::RegisterKernelWrapper( const char *kernelName, cl_kernel_function function )
+int OpenclDevice::registerKernelWrapper( const char *kernelName, cl_kernel_function function )
{
int i;
- //printf("oclwrapper:RegisterKernelWrapper...%d\n", gpuEnv.mnKernelCount);
+ //printf("oclwrapper:registerKernelWrapper...%d\n", gpuEnv.mnKernelCount);
for ( i = 0; i < gpuEnv.mnKernelCount; i++ )
{
if ( strcasecmp( kernelName, gpuEnv.mArrykernelNames[i]) == 0 )
@@ -892,13 +928,13 @@ int OpenclDevice::RegisterKernelWrapper( const char *kernelName, cl_kernel_funct
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;
}
@@ -908,54 +944,46 @@ OclCalc::OclCalc()
fprintf(stderr,"OclCalc:: init opencl ...\n");
nFormulaColSize = 0;
nFormulaRowSize = 0;
+ nArithmeticLen = 0;
+ nFormulaLen = 0;
+ mpClmemSrcData = NULL;
+ mpClmemStartPos = NULL;
+ mpClmemEndPos = NULL;
+ mpClmemLeftData = NULL;
+ mpClmemRightData = NULL;
+ mpClmemMergeLfData = NULL;
+ mpClmemMergeRtData = NULL;
+ mpClmemMatixSumSize = NULL;
+ mpClmemeOp = NULL;
}
OclCalc::~OclCalc()
{
- fprintf(stderr,"OclCalc:: opencl end ...\n");
+ releaseOclBuffer();
}
-
-/////////////////////////////////////////////////////////////////////////////
-int OclCalc::CreateBuffer64Bits( double *&dpSrcData, uint *&npStartPos, uint *&npEndPos, int nBufferSize )
+int OclCalc::releaseOclBuffer(void)
{
cl_int clStatus = 0;
- SetKernelEnv( &kEnv );
-
- mpClmemSrcData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
- nBufferSize * sizeof(double), NULL, &clStatus );
- CHECK_OPENCL( clStatus, "clCreateBuffer" );
- mpClmemStartPos = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
- nBufferSize * sizeof(unsigned int), NULL, &clStatus );
- CHECK_OPENCL( clStatus, "clCreateBuffer" );
- mpClmemEndPos = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
- nBufferSize * sizeof(unsigned int), NULL, &clStatus );
- CHECK_OPENCL( clStatus, "clCreateBuffer" );
-
- dpSrcData = (double *)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmemSrcData, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, 0,
- nBufferSize * sizeof(double), 0, NULL, NULL, &clStatus );
- CHECK_OPENCL(clStatus,"clEnqueueMapBuffer");
- clFinish(kEnv.mpkCmdQueue);
- npStartPos = (uint *)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmemStartPos, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, 0,
- nBufferSize * sizeof(uint), 0, NULL, NULL, &clStatus );
- CHECK_OPENCL( clStatus, "clEnqueueMapBuffer" );
- clFinish( kEnv.mpkCmdQueue );
- npEndPos = (uint *)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmemEndPos, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, 0,
- nBufferSize * sizeof(uint), 0, NULL, NULL, &clStatus );
- CHECK_OPENCL( clStatus,"clEnqueueMapBuffer" );
- clFinish( kEnv.mpkCmdQueue );
- //printf("In CreateBuffer, pptrr is %d,%d,%d\n",fpSrcData,npStartPos,npEndPos);
- return 0;
+ CHECK_OPENCL_RELEASE( clStatus, mpClmemSrcData );
+ CHECK_OPENCL_RELEASE( clStatus, mpClmemStartPos );
+ CHECK_OPENCL_RELEASE( clStatus, mpClmemEndPos );
+ CHECK_OPENCL_RELEASE( clStatus, mpClmemLeftData );
+ CHECK_OPENCL_RELEASE( clStatus, mpClmemRightData );
+ fprintf(stderr,"OclCalc:: opencl end ...\n");
+ return 1;
}
-int OclCalc::CreateBuffer64Bits( double *&dpLeftData, double *&dpRightData, int nBufferSize )
+/////////////////////////////////////////////////////////////////////////////
+
+int OclCalc::createBuffer64Bits( double *&dpLeftData, double *&dpRightData, int nBufferSize )
{
cl_int clStatus = 0;
- SetKernelEnv( &kEnv );
+ setKernelEnv( &kEnv );
- mpClmemLeftData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
+ mpClmemLeftData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR ),
nBufferSize * sizeof(double), NULL, &clStatus );
CHECK_OPENCL( clStatus, "clCreateBuffer" );
- mpClmemRightData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
+ mpClmemRightData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR ),
nBufferSize * sizeof(double), NULL, &clStatus );
CHECK_OPENCL( clStatus,"clCreateBuffer" );
dpLeftData = (double *)clEnqueueMapBuffer( kEnv.mpkCmdQueue,mpClmemLeftData,CL_TRUE,CL_MAP_WRITE,0,
@@ -969,54 +997,169 @@ int OclCalc::CreateBuffer64Bits( double *&dpLeftData, double *&dpRightData, int
//printf("In CreateBuffer, pptrr is %d,%d,%d\n",fpSrcData,npStartPos,npEndPos);
return 0;
}
-int OclCalc::CreateBuffer64Bits( double *&dpSumProMergeLfData, double *&dpSumProMergeRtData, uint *&npSumSize, int nMatixSize, int nBufferSize )
+
+int OclCalc::mapAndCopy64Bits(const double *dpTempSrcData,unsigned int *unStartPos,unsigned int *unEndPos,int nBufferSize ,int nRowsize)
{
cl_int clStatus = 0;
- SetKernelEnv( &kEnv );
- mpClmemMergeLfData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
- nMatixSize * sizeof(double), NULL, &clStatus );
- CHECK_OPENCL( clStatus, "clCreateBuffer" );
- mpClmemMergeRtData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
- nMatixSize * sizeof(double), NULL, &clStatus );
- CHECK_OPENCL( clStatus, "clCreateBuffer" );
- mpClmemMatixSumSize = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR),
- nMatixSize * sizeof(unsigned int), NULL, &clStatus );
- CHECK_OPENCL( clStatus, "clCreateBuffer" );
+ double * dpSrcDataMap = (double *)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmemSrcData, CL_TRUE,CL_MAP_WRITE, 0,
+ nBufferSize * sizeof(double), 0, NULL, NULL, &clStatus );
+ CHECK_OPENCL( clStatus, "clEnqueueMapBuffer" );
+ clFinish( kEnv.mpkCmdQueue );
+ unsigned int *npStartPosMap = (uint *)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmemStartPos, CL_TRUE,CL_MAP_WRITE, 0,
+ nRowsize * sizeof(uint), 0, NULL, NULL, &clStatus );
+ CHECK_OPENCL( clStatus, "clEnqueueMapBuffer" );
+ clFinish( kEnv.mpkCmdQueue );
+ unsigned int *npEndPosMap = (uint *)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmemEndPos, CL_TRUE, CL_MAP_WRITE, 0,
+ nRowsize * sizeof(uint), 0, NULL, NULL, &clStatus );
+ CHECK_OPENCL( clStatus, "clEnqueueMapBuffer" );
+ clFinish( kEnv.mpkCmdQueue );
+ for(int i=0;i<nBufferSize;i++)
+ dpSrcDataMap[i] = dpTempSrcData[i];
+ for(int i=0;i<nRowsize;i++)
+ {
+ npStartPosMap[i] = unStartPos[i];
+ npEndPosMap[i] = unEndPos[i];
+ }
+ clStatus = clEnqueueUnmapMemObject( kEnv.mpkCmdQueue, mpClmemSrcData, dpSrcDataMap, 0, NULL, NULL );
+ CHECK_OPENCL( clStatus, "clEnqueueUnmapMemObject" );
+ clFinish( kEnv.mpkCmdQueue );
... etc. - the rest is truncated
More information about the Libreoffice-commits
mailing list