[Libreoffice-commits] core.git: Branch 'feature/calc-group-interpreter-3' - 75 commits - config_host.mk.in configure.ac officecfg/registry RepositoryExternal.mk Repository.mk RepositoryModule_host.mk sc/CppunitTest_sc_ucalc.mk sc/inc sc/Library_scfilt.mk sc/Library_sc.mk sc/Library_scopencl.mk sc/Module_sc.mk sc/source sc/uiconfig sc/workben
Markus Mohrhard
markus.mohrhard at googlemail.com
Thu Sep 19 08:11:09 PDT 2013
Repository.mk | 1
RepositoryExternal.mk | 17
RepositoryModule_host.mk | 2
config_host.mk.in | 2
configure.ac | 52
officecfg/registry/schema/org/openoffice/Office/Calc.xcs | 14
sc/CppunitTest_sc_ucalc.mk | 3
sc/Library_sc.mk | 10
sc/Library_scfilt.mk | 2
sc/Library_scopencl.mk | 48
sc/Module_sc.mk | 6
sc/inc/calcconfig.hxx | 4
sc/inc/document.hxx | 2
sc/inc/formulagroup.hxx | 8
sc/inc/pch/precompiled_scopencl.cxx | 12
sc/inc/pch/precompiled_scopencl.hxx | 31
sc/inc/platforminfo.hxx | 45
sc/source/core/opencl/clcc/clew.cxx | 326 ++
sc/source/core/opencl/clcc/clew.h | 1318 +++++++++++
sc/source/core/opencl/formulagroupcl.cxx | 184 -
sc/source/core/opencl/oclkernels.hxx | 5
sc/source/core/opencl/openclwrapper.cxx | 1700 ++++++++-------
sc/source/core/opencl/openclwrapper.hxx | 287 +-
sc/source/core/tool/calcconfig.cxx | 7
sc/source/core/tool/formulagroup.cxx | 168 +
sc/source/core/tool/formulaopt.cxx | 42
sc/source/core/tool/platforminfo.cxx | 21
sc/source/filter/excel/excform.cxx | 92
sc/source/filter/excel/impop.cxx | 83
sc/source/filter/excel/read.cxx | 16
sc/source/filter/ftools/clkernelthread.cxx | 28
sc/source/filter/inc/clkernelthread.hxx | 26
sc/source/filter/inc/excform.hxx | 4
sc/source/filter/inc/imp_op.hxx | 31
sc/source/filter/inc/workbookhelper.hxx | 2
sc/source/filter/oox/workbookhelper.cxx | 23
sc/source/filter/oox/worksheetfragment.cxx | 9
sc/source/ui/app/scmod.cxx | 3
sc/source/ui/optdlg/calcoptionsdlg.cxx | 140 +
sc/source/ui/optdlg/calcoptionsdlg.hxx | 26
sc/uiconfig/scalc/ui/formulacalculationoptions.ui | 211 +
sc/workben/opencl/platform_detect.cxx | 72
42 files changed, 3906 insertions(+), 1177 deletions(-)
New commits:
commit b3a68951367a6aa2e4ad7a6114dd6312c0f85f40
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Thu Sep 19 17:02:10 2013 +0200
fix windows build with pch
diff --git a/sc/inc/pch/precompiled_scopencl.cxx b/sc/inc/pch/precompiled_scopencl.cxx
new file mode 100644
index 0000000..85d15b4
--- /dev/null
+++ b/sc/inc/pch/precompiled_scopencl.cxx
@@ -0,0 +1,12 @@
+/* -*- 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 "precompiled_scopencl.hxx"
+
+/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/inc/pch/precompiled_scopencl.hxx b/sc/inc/pch/precompiled_scopencl.hxx
new file mode 100644
index 0000000..b4cc3cf
--- /dev/null
+++ b/sc/inc/pch/precompiled_scopencl.hxx
@@ -0,0 +1,31 @@
+/* -*- 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/.
+ */
+
+/*
+ This file has been autogenerated by update_pch.sh . It is possible to edit it
+ manually (such as when an include file has been moved/renamed/removed. All such
+ manual changes will be rewritten by the next run of update_pch.sh (which presumably
+ also fixes all possible problems, so it's usually better to use it).
+*/
+
+#include "formula/vectortoken.hxx"
+#include "sal/config.h"
+#include <boost/scoped_array.hpp>
+#include <cmath>
+#include <config_folders.h>
+#include <osl/file.hxx>
+#include <rtl/bootstrap.hxx>
+#include <rtl/digest.h>
+#include <rtl/strbuf.hxx>
+#include <rtl/ustring.hxx>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
commit c74018fcc7436d57a0f2d53036ae16e6d7e2e662
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Wed Sep 18 18:09:02 2013 -0400
"what!?" - Let's not say that...
Change-Id: I1c9c2a62b2a1ede08fb2f27f541b501dc3ce7b2f
diff --git a/sc/source/filter/excel/excform.cxx b/sc/source/filter/excel/excform.cxx
index 4c49461..0c34957 100644
--- a/sc/source/filter/excel/excform.cxx
+++ b/sc/source/filter/excel/excform.cxx
@@ -142,8 +142,6 @@ void ImportExcel::Formula(
GetXFRangeBuffer().SetXF(aScPos, nXF);
SetLastFormula(aScPos.Col(), aScPos.Row(), fCurVal, nXF, pCell);
}
- else
- fprintf(stdout, "ImportExcel::Formula: what!?\n");
}
else
{
commit b92e53513bb29f988c80577e36cf135547bc940f
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Wed Sep 18 17:40:31 2013 -0400
Work around another Excel bug with incorrect shared formula range.
Take the start row position from the preceding formula record, rather
than believing what's in the shared formula record. The latter can be
wrong, and can be wrong often.
Change-Id: I3a4da110727a7719e5f8cb3e6250c0e1bef04c64
diff --git a/sc/source/filter/excel/excform.cxx b/sc/source/filter/excel/excform.cxx
index b557de5..4c49461 100644
--- a/sc/source/filter/excel/excform.cxx
+++ b/sc/source/filter/excel/excform.cxx
@@ -102,6 +102,9 @@ void ImportExcel::Formula4()
void ImportExcel::Formula(
const XclAddress& rXclPos, sal_uInt16 nXF, sal_uInt16 nFormLen, double fCurVal, bool bShrFmla)
{
+ if (!nFormLen)
+ return;
+
ScAddress aScPos( ScAddress::UNINITIALIZED );
if (!GetAddressConverter().ConvertAddress(aScPos, rXclPos, GetCurrScTab(), true))
// Conversion failed.
@@ -116,22 +119,39 @@ void ImportExcel::Formula(
if (bShrFmla)
{
// This is a shared formula. Get the token array from the shared formula pool.
- ScFormulaCellGroupRef xGroup = pFormConv->GetSharedFormula(maStrm, aScPos.Col(), nFormLen);
- if (xGroup)
+ SCCOL nSharedCol;
+ SCROW nSharedRow;
+ if (pFormConv->ReadSharedFormulaPosition(maStrm, nSharedCol, nSharedRow))
{
- if (xGroup->mnStart == aScPos.Row())
- // Generate code for the top cell only.
- xGroup->compileCode(rDoc.getDoc(), aScPos, formula::FormulaGrammar::GRAM_DEFAULT);
-
- ScFormulaCell* pCell = new ScFormulaCell(pD, aScPos, xGroup);
- rDoc.getDoc().EnsureTable(aScPos.Tab());
- rDoc.setFormulaCell(aScPos, pCell);
- xGroup->mnLength = aScPos.Row() - xGroup->mnStart + 1;
- pCell->SetNeedNumberFormat(false);
- if (!rtl::math::isNan(fCurVal))
- pCell->SetResultDouble(fCurVal);
-
- GetXFRangeBuffer().SetXF(aScPos, nXF);
+ ScAddress aRefPos(aScPos.Col(), nSharedRow, GetCurrScTab());
+ ScFormulaCellGroupRef xGroup = pFormConv->GetSharedFormula(aRefPos);
+ if (xGroup)
+ {
+ // Make sure the this one follows immediately below another shared formula cell.
+ LastFormula* pLast = GetLastFormula(aScPos.Col());
+ if (pLast && pLast->mpCell && pLast->mnRow == (aScPos.Row()-1))
+ {
+ ScFormulaCell* pCell = new ScFormulaCell(pD, aScPos, xGroup);
+ rDoc.getDoc().EnsureTable(aScPos.Tab());
+ rDoc.setFormulaCell(aScPos, pCell);
+ xGroup->mnLength = aScPos.Row() - xGroup->mnStart + 1;
+ pCell->SetNeedNumberFormat(false);
+ if (!rtl::math::isNan(fCurVal))
+ pCell->SetResultDouble(fCurVal);
+
+ GetXFRangeBuffer().SetXF(aScPos, nXF);
+ SetLastFormula(aScPos.Col(), aScPos.Row(), fCurVal, nXF, pCell);
+ }
+ else
+ fprintf(stdout, "ImportExcel::Formula: what!?\n");
+ }
+ else
+ {
+ // Shared formula not found even though it's clearly a shared formula.
+ // The cell will be created in the following shared formula
+ // record.
+ SetLastFormula(aScPos.Col(), aScPos.Row(), fCurVal, nXF, NULL);
+ }
return;
}
}
@@ -158,6 +178,7 @@ void ImportExcel::Formula(
pCell = new ScFormulaCell(&rDoc.getDoc(), aScPos, pResult);
rDoc.getDoc().EnsureTable(aScPos.Tab());
rDoc.setFormulaCell(aScPos, pCell);
+ SetLastFormula(aScPos.Col(), aScPos.Row(), fCurVal, nXF, pCell);
}
else
{
@@ -1692,30 +1713,30 @@ const ScTokenArray* ExcelToSc::GetBoolErr( XclBoolError eType )
return pErgebnis;
}
-ScFormulaCellGroupRef ExcelToSc::GetSharedFormula( XclImpStream& aIn, SCCOL nCol, sal_Size nFormulaLen )
+bool ExcelToSc::ReadSharedFormulaPosition( XclImpStream& rStrm, SCCOL& rCol, SCROW& rRow )
{
- if (!nFormulaLen)
- return ScFormulaCellGroupRef();
-
- aIn.PushPosition();
+ rStrm.PushPosition();
sal_uInt8 nOp;
- aIn >> nOp;
+ rStrm >> nOp;
if (nOp != 0x01) // must be PtgExp token.
{
- aIn.PopPosition();
- return ScFormulaCellGroupRef();
+ rStrm.PopPosition();
+ return false;
}
- sal_uInt16 nLeftCol, nRow;
- aIn >> nRow >> nLeftCol;
-
- ScAddress aRefPos(nCol, nRow, GetCurrScTab());
- ScFormulaCellGroupRef xGroup = GetOldRoot().pShrfmlaBuff->Find(aRefPos);
+ sal_uInt16 nRow, nCol;
+ rStrm >> nRow >> nCol;
+ rStrm.PopPosition();
+ rCol = nCol;
+ rRow = nRow;
+ return true;
+}
- aIn.PopPosition();
- aIn.Ignore(nFormulaLen);
+ScFormulaCellGroupRef ExcelToSc::GetSharedFormula( const ScAddress& rRefPos )
+{
+ ScFormulaCellGroupRef xGroup = GetOldRoot().pShrfmlaBuff->Find(rRefPos);
return xGroup;
}
diff --git a/sc/source/filter/excel/impop.cxx b/sc/source/filter/excel/impop.cxx
index 16c4ac5..a06aac1 100644
--- a/sc/source/filter/excel/impop.cxx
+++ b/sc/source/filter/excel/impop.cxx
@@ -114,8 +114,10 @@ ImportExcel::ImportExcel( XclImpRootData& rImpData, SvStream& rStrm ):
maStrm( rStrm, GetRoot() ),
aIn( maStrm ),
maScOleSize( ScAddress::INITIALIZE_INVALID ),
+ mpLastFormula(NULL),
mnLastRefIdx( 0 ),
mnIxfeIndex( 0 ),
+ mnLastRecId(0),
mbBiff2HasXfs(false),
mbBiff2HasXfsValid(false),
mbRunCLKernelThread(true)
@@ -165,6 +167,31 @@ ImportExcel::~ImportExcel( void )
delete pFormConv;
}
+void ImportExcel::SetLastFormula( SCCOL nCol, SCROW nRow, double fVal, sal_uInt16 nXF, ScFormulaCell* pCell )
+{
+ LastFormulaMapType::iterator it = maLastFormulaCells.find(nCol);
+ if (it == maLastFormulaCells.end())
+ {
+ std::pair<LastFormulaMapType::iterator, bool> r =
+ maLastFormulaCells.insert(
+ LastFormulaMapType::value_type(nCol, LastFormula()));
+ it = r.first;
+ }
+
+ it->second.mnCol = nCol;
+ it->second.mnRow = nRow;
+ it->second.mpCell = pCell;
+ it->second.mfValue = fVal;
+ it->second.mnXF = nXF;
+
+ mpLastFormula = &it->second;
+}
+
+ImportExcel::LastFormula* ImportExcel::GetLastFormula( SCCOL nCol )
+{
+ LastFormulaMapType::iterator it = maLastFormulaCells.find(nCol);
+ return it == maLastFormulaCells.end() ? NULL : &it->second;
+}
void ImportExcel::ReadFileSharing()
{
@@ -815,6 +842,21 @@ void ImportExcel::Standardwidth( void )
void ImportExcel::Shrfmla( void )
{
+ switch (mnLastRecId)
+ {
+ case EXC_ID2_FORMULA:
+ case EXC_ID3_FORMULA:
+ case EXC_ID4_FORMULA:
+ // This record MUST immediately follow a FORMULA record.
+ break;
+ default:
+ return;
+ }
+
+ if (!mpLastFormula)
+ // The last FORMULA record should have left this data.
+ return;
+
sal_uInt16 nFirstRow, nLastRow, nLenExpr;
sal_uInt8 nFirstCol, nLastCol;
@@ -829,13 +871,37 @@ void ImportExcel::Shrfmla( void )
pFormConv->Reset();
pFormConv->Convert( pErgebnis, maStrm, nLenExpr, true, FT_SharedFormula );
-
OSL_ENSURE( pErgebnis, "+ImportExcel::Shrfmla(): ScTokenArray is NULL!" );
- pExcRoot->pShrfmlaBuff->Store( ScRange( static_cast<SCCOL>(nFirstCol),
- static_cast<SCROW>(nFirstRow), GetCurrScTab(),
- static_cast<SCCOL>(nLastCol), static_cast<SCROW>(nLastRow),
- GetCurrScTab()), *pErgebnis );
+ // The range in this record can be erroneous especially the row range.
+ // Use the row from the last FORMULA record as the start row. The end row
+ // will be adjusted by the formula cells that follow.
+ SCCOL nCol1 = nFirstCol;
+ SCCOL nCol2 = nLastCol;
+ SCROW nRow1 = mpLastFormula->mnRow;
+
+ pExcRoot->pShrfmlaBuff->Store(
+ ScRange(nCol1, nRow1, GetCurrScTab(), nCol2, nRow1, GetCurrScTab()), *pErgebnis);
+
+ // Create formula cell for the last formula record.
+
+ ScAddress aPos(nCol1, nRow1, GetCurrScTab());
+ ScFormulaCellGroupRef xGroup = pExcRoot->pShrfmlaBuff->Find(aPos);
+ if (xGroup)
+ {
+ ScDocumentImport& rDoc = GetDocImport();
+ xGroup->compileCode(rDoc.getDoc(), aPos, formula::FormulaGrammar::GRAM_DEFAULT);
+
+ ScFormulaCell* pCell = new ScFormulaCell(pD, aPos, xGroup);
+ rDoc.getDoc().EnsureTable(aPos.Tab());
+ rDoc.setFormulaCell(aPos, pCell);
+ pCell->SetNeedNumberFormat(false);
+ if (!rtl::math::isNan(mpLastFormula->mfValue))
+ pCell->SetResultDouble(mpLastFormula->mfValue);
+
+ GetXFRangeBuffer().SetXF(aPos, mpLastFormula->mnXF);
+ mpLastFormula->mpCell = pCell;
+ }
}
@@ -1181,6 +1247,8 @@ void ImportExcel::NeueTabelle( void )
}
pExcRoot->pShrfmlaBuff->Clear();
+ maLastFormulaCells.clear();
+ mpLastFormula = NULL;
InitializeTable( nTab );
diff --git a/sc/source/filter/excel/read.cxx b/sc/source/filter/excel/read.cxx
index 14396f6..610db55 100644
--- a/sc/source/filter/excel/read.cxx
+++ b/sc/source/filter/excel/read.cxx
@@ -630,7 +630,6 @@ FltError ImportExcel::Read( void )
case 0x0223: Externname34(); break; // EXTERNNAME [ 34 ]
case 0x0225: Defrowheight345();break;//DEFAULTROWHEI[ 345]
case 0x023E: rTabViewSett.ReadWindow2( maStrm, false );break;
- case 0x04BC: Shrfmla(); break; // SHRFMLA [ 5]
}
}
}
@@ -654,6 +653,7 @@ FltError ImportExcel::Read( void )
case EXC_ID2_FORMULA:
case EXC_ID3_FORMULA:
case EXC_ID4_FORMULA: Formula25(); break;
+ case EXC_ID_SHRFMLA: Shrfmla(); break;
case 0x0A: Eof(); eAkt = Z_Biff5E; break;
case 0x14:
case 0x15: rPageSett.ReadHeaderFooter( maStrm ); break;
@@ -833,7 +833,9 @@ FltError ImportExcel8::Read( void )
std::vector<OUString> aCodeNames;
std::vector < SCTAB > nTabsWithNoCodeName;
- while( eAkt != EXC_STATE_END )
+ sal_uInt16 nRecId = 0;
+
+ for (; eAkt != EXC_STATE_END; mnLastRecId = nRecId)
{
if( eAkt == EXC_STATE_BEFORE_SHEET )
{
@@ -937,7 +939,7 @@ FltError ImportExcel8::Read( void )
if( eAkt != EXC_STATE_SHEET_PRE && eAkt != EXC_STATE_GLOBALS_PRE )
pProgress->ProgressAbs( nProgressBaseSize + aIn.GetSvStreamPos() - nProgressBasePos );
- sal_uInt16 nRecId = aIn.GetRecId();
+ nRecId = aIn.GetRecId();
/* #i39464# Ignore records between USERSVIEWBEGIN and USERSVIEWEND
completely (user specific view settings). Otherwise view settings
@@ -1145,7 +1147,6 @@ FltError ImportExcel8::Read( void )
case EXC_ID2_ARRAY:
case EXC_ID3_ARRAY: Array34(); break; // ARRAY [ 34 ]
case 0x0225: Defrowheight345();break;//DEFAULTROWHEI[ 345 ]
- case EXC_ID_SHRFMLA: Shrfmla(); break; // SHRFMLA [ 5 ]
case 0x0867: SheetProtection(); break; // SHEETPROTECTION
}
}
@@ -1179,6 +1180,7 @@ FltError ImportExcel8::Read( void )
case EXC_ID2_FORMULA:
case EXC_ID3_FORMULA:
case EXC_ID4_FORMULA: Formula25(); break;
+ case EXC_ID_SHRFMLA: Shrfmla(); break;
case 0x000C: Calccount(); break; // CALCCOUNT
case 0x0010: Delta(); break; // DELTA
case 0x0011: Iteration(); break; // ITERATION
diff --git a/sc/source/filter/inc/excform.hxx b/sc/source/filter/inc/excform.hxx
index 0e2004c..11254a0 100644
--- a/sc/source/filter/inc/excform.hxx
+++ b/sc/source/filter/inc/excform.hxx
@@ -63,7 +63,9 @@ public:
void GetDummy( const ScTokenArray*& );
const ScTokenArray* GetBoolErr( XclBoolError );
- ScFormulaCellGroupRef GetSharedFormula( XclImpStream& rStrm, SCCOL nCol, sal_Size nFormulaLen );
+
+ bool ReadSharedFormulaPosition( XclImpStream& rStrm, SCCOL& rCol, SCROW& rRow );
+ ScFormulaCellGroupRef GetSharedFormula( const ScAddress& rRefPos );
static void SetError( ScFormulaCell& rCell, const ConvErr eErr );
diff --git a/sc/source/filter/inc/imp_op.hxx b/sc/source/filter/inc/imp_op.hxx
index 6e44f31..7c58395 100644
--- a/sc/source/filter/inc/imp_op.hxx
+++ b/sc/source/filter/inc/imp_op.hxx
@@ -35,7 +35,7 @@
#include <boost/shared_ptr.hpp>
#include <boost/ptr_container/ptr_vector.hpp>
-
+#include <boost/unordered_map.hpp>
class SvStream;
@@ -82,6 +82,16 @@ private:
class ImportExcel : public ImportTyp, protected XclImpRoot
{
protected:
+ struct LastFormula
+ {
+ SCCOL mnCol;
+ SCROW mnRow;
+ double mfValue;
+ sal_uInt16 mnXF;
+ ScFormulaCell* mpCell;
+ };
+ typedef boost::unordered_map<SCCOL, LastFormula> LastFormulaMapType;
+
rtl::Reference<sc::CLBuildKernelThread> mxCLKernelThread;
static const double fExcToTwips; // Umrechnung 1/256 Zeichen -> Twips
@@ -104,11 +114,14 @@ protected:
typedef boost::ptr_vector< XclImpOutlineDataBuffer > XclImpOutlineListBuffer;
XclImpOutlineListBuffer* pOutlineListBuffer;
+ LastFormulaMapType maLastFormulaCells; // Keep track of last formula cells in each column.
+ LastFormula* mpLastFormula;
+
sal_Int16 mnLastRefIdx;
sal_uInt16 mnIxfeIndex; /// Current XF identifier from IXFE record.
+ sal_uInt16 mnLastRecId;
SCTAB nBdshtTab; // Counter fuer Boundsheet
- ScFormulaCell* pLastFormCell; // fuer String-Records
sal_Bool bTabTruncated; // wenn Bereichsueberschreitung zum
// Abschneiden von Zellen fuehrt
@@ -117,6 +130,9 @@ protected:
bool mbBiff2HasXfsValid:1; /// False = mbBiff2HasXfs is undetermined yet.
bool mbRunCLKernelThread:1;
+ void SetLastFormula( SCCOL nCol, SCROW nRow, double fVal, sal_uInt16 nXF, ScFormulaCell* pCell );
+ LastFormula* GetLastFormula( SCCOL nCol );
+
// Record-Funktionen
void ReadFileSharing();
commit 36a199438ba96797171d907ad423357cc9cedcfe
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Wed Sep 18 16:47:15 2013 +0200
fix storing of open cl device in configuration files
Change-Id: Icc343b55fe40c430b87d2e4c98701d78c61e3a96
diff --git a/officecfg/registry/schema/org/openoffice/Office/Calc.xcs b/officecfg/registry/schema/org/openoffice/Office/Calc.xcs
index 25fe724..4c583a4 100644
--- a/officecfg/registry/schema/org/openoffice/Office/Calc.xcs
+++ b/officecfg/registry/schema/org/openoffice/Office/Calc.xcs
@@ -1344,12 +1344,12 @@
</info>
<value>true</value>
</prop>
- <prop oor:name="OpenCLAutoDevice" oor:type="xs:string" oor:nillable="false">
+ <prop oor:name="OpenCLDevice" oor:type="xs:string" oor:nillable="false">
<!-- UIHints: Tools - Options Spreadsheet Formula -->
<info>
<desc>The Device ID of the OpenCL device selected if OpenCLAutoSelect is false</desc>
</info>
- <value></value>
+ <value/>
</prop>
</group>
<group oor:name="Syntax">
diff --git a/sc/source/core/tool/formulaopt.cxx b/sc/source/core/tool/formulaopt.cxx
index f649927..1ce5e09 100644
--- a/sc/source/core/tool/formulaopt.cxx
+++ b/sc/source/core/tool/formulaopt.cxx
@@ -216,7 +216,7 @@ Sequence<OUString> ScFormulaCfg::GetPropertyNames()
"Load/ODFRecalcMode", // SCFORMULAOPT_ODF_RECALC
"Calculation/OpenCL", // SCFORMULAOPT_OPENCL_ENABLED
"Calculation/OpenCLAutoSelect", // SCFORMULAOPT_OPENCL_AUTOSELECT
- "Calculation/OpenCLAutoDevice" // SCFORMULAOPT_OPENCL_DEVICE
+ "Calculation/OpenCLDevice" // SCFORMULAOPT_OPENCL_DEVICE
};
Sequence<OUString> aNames(SCFORMULAOPT_COUNT);
OUString* pNames = aNames.getArray();
@@ -428,6 +428,7 @@ void ScFormulaCfg::UpdateFromProperties( const Sequence<OUString>& aNames )
pValues[nProp] >>= aOpenCLDevice;
GetCalcConfig().maOpenCLDevice = aOpenCLDevice;
}
+ break;
default:
;
}
@@ -540,6 +541,8 @@ void ScFormulaCfg::Commit()
{
sal_Bool bVal = GetCalcConfig().mbOpenCLAutoSelect;
pValues[nProp] <<= bVal;
+ sc::FormulaGroupInterpreter::switchOpenCLDevice(
+ GetCalcConfig().maOpenCLDevice, bVal);
}
break;
case SCFORMULAOPT_OPENCL_DEVICE:
commit 7facca58d55d1afe2c3e90b58bf00304293ec698
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Wed Sep 18 16:26:08 2013 +0200
small fixes for broken code
Change-Id: I77afc818d5575523077328042ac9418f521edc9c
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index d4cf1b7..d273463 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -251,8 +251,8 @@ std::vector<boost::shared_ptr<osl::File> > OpenclDevice::binaryGenerated( const
}
}
}
- return aGeneratedFiles;
+ return aGeneratedFiles;
}
int OpenclDevice::writeBinaryToFile( const OString& rFileName, const char* binary, size_t numBytes )
@@ -270,7 +270,6 @@ int OpenclDevice::writeBinaryToFile( const OString& rFileName, const char* binar
assert(numBytes == nBytesWritten);
return 1;
-
}
int OpenclDevice::generatBinFromKernelSource( cl_program program, const char * clFileName )
@@ -414,9 +413,7 @@ int OpenclDevice::cachedOfKernerPrg( const GPUEnv *gpuEnvCached, const char * cl
int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
{
cl_int clStatus = 0;
- size_t length;
int binary_status, idx;
- cl_uint numDevices;
const char* filename = "kernel.cl";
fprintf(stderr, "compileKernelFile ... \n");
if ( cachedOfKernerPrg(gpuInfo, filename) == 1 )
@@ -426,6 +423,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
idx = gpuInfo->mnFileCount;
+ cl_uint numDevices;
clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES,
sizeof(numDevices), &numDevices, NULL );
CHECK_OPENCL( clStatus, "clGetContextInfo" );
@@ -437,6 +435,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
if (aGeneratedFiles.size() == numDevices)
{
bBinaryExisted = true;
+ boost::scoped_array<size_t> length(new size_t[numDevices]);
boost::scoped_array<unsigned char*> pBinary(new unsigned char*[numDevices]);
for(size_t i = 0; i < numDevices; ++i)
{
@@ -448,7 +447,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
if(nSize != nBytesRead)
assert(false);
- length = nBytesRead;
+ length[i] = nBytesRead;
pBinary[i] = binary;
}
@@ -469,7 +468,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
fprintf(stderr, "Create kernel from binary\n");
gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary( gpuInfo->mpContext,numDevices,
- mpArryDevsID.get(), &length, (const unsigned char**) pBinary.get(),
+ mpArryDevsID.get(), length.get(), (const unsigned char**) pBinary.get(),
&binary_status, &clStatus );
if(clStatus != CL_SUCCESS)
{
@@ -515,6 +514,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
if ( clStatus != CL_SUCCESS )
{
+ size_t length;
printf ("BuildProgram error!\n");
if ( !gpuInfo->mnIsUserCreated )
{
commit e96d8c9e9c8a6fde50937bcb9b54fddf9dc7fc55
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Wed Sep 18 16:21:43 2013 +0200
select maximum one opencl device
Change-Id: If06caaa6b39131359ffc81fadad200c8cf8fe24a
diff --git a/sc/source/ui/optdlg/calcoptionsdlg.cxx b/sc/source/ui/optdlg/calcoptionsdlg.cxx
index 9607fb6..eefb7f6 100644
--- a/sc/source/ui/optdlg/calcoptionsdlg.cxx
+++ b/sc/source/ui/optdlg/calcoptionsdlg.cxx
@@ -204,6 +204,8 @@ void ScCalcOptionsDialog::fillOpenclList()
OUString aStoredDevice = maConfig.maOpenCLDevice;
+ SvTreeListEntry* pSelectedEntry = NULL;
+
sc::FormulaGroupInterpreter::fillOpenCLInfo(maPlatformInfo);
for(std::vector<sc::OpenclPlatformInfo>::iterator it = maPlatformInfo.begin(),
itEnd = maPlatformInfo.end(); it != itEnd; ++it)
@@ -215,13 +217,18 @@ void ScCalcOptionsDialog::fillOpenclList()
SvTreeListEntry* pEntry = mpOpenclInfoList->InsertEntry(aDeviceId);
if(aDeviceId == aStoredDevice)
{
- mpOpenclInfoList->GetModel()->GetView(0)->Select(pEntry);
+ pSelectedEntry = pEntry;
}
pEntry->SetUserData(&(*itr));
}
}
mpOpenclInfoList->SetUpdateMode(true);
+ mpOpenclInfoList->GetModel()->GetView(0)->SelectAll(false, false);
+ if(pSelectedEntry)
+ {
+ mpOpenclInfoList->GetModel()->GetView(0)->Select(pSelectedEntry);
+ }
SelectedDeviceChanged();
}
commit 5c31bc44ada3ae36c27c550d33805c53a0bdcb0f
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Wed Sep 18 13:06:16 2013 +0200
fix memory leak
Change-Id: If1afe59bc5bb40bb3fff4c74b863c6d77e006cc5
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 0fe2220..d4cf1b7 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -475,10 +475,10 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
{
// something went wrong, fall back to compiling from source
bBinaryExisted = false;
- for(size_t i = 0; i < numDevices; ++i)
- {
- delete[] pBinary[i];
- }
+ }
+ for(size_t i = 0; i < numDevices; ++i)
+ {
+ delete[] pBinary[i];
}
}
commit 852e88a0ae2df662d001a6ec67d190160a1423a2
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Wed Sep 18 13:05:31 2013 +0200
fall back to compiling from source if binary failed
Change-Id: I8e8d8fb5b1da18c44e1537f97317d4e5b57c47b6
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index f469aac..0fe2220 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -473,14 +473,16 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
&binary_status, &clStatus );
if(clStatus != CL_SUCCESS)
{
+ // something went wrong, fall back to compiling from source
+ bBinaryExisted = false;
for(size_t i = 0; i < numDevices; ++i)
{
delete[] pBinary[i];
}
}
- CHECK_OPENCL( clStatus, "clCreateProgramWithBinary" );
}
- else
+
+ if(!bBinaryExisted)
{
// create a CL program using the kernel source
fprintf(stderr, "Create kernel from source\n");
commit 2a9b91dcca904304e6ebd6ec8566f1acce9d51f9
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Wed Sep 18 12:45:37 2013 +0200
void* is dangerous
no idea how that ever worked
Change-Id: Ie14c86a241a2fa7e1963385987245cd0ee965fcf
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 1a19451..f469aac 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -437,17 +437,19 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
if (aGeneratedFiles.size() == numDevices)
{
bBinaryExisted = true;
- boost::scoped_array<char*> pBinary(new char*[numDevices]);
+ boost::scoped_array<unsigned char*> pBinary(new unsigned char*[numDevices]);
for(size_t i = 0; i < numDevices; ++i)
{
sal_uInt64 nSize;
aGeneratedFiles[i]->getSize(nSize);
- char* binary = new char[nSize];
+ unsigned char* binary = new unsigned char[nSize];
sal_uInt64 nBytesRead;
aGeneratedFiles[i]->read(binary, nSize, nBytesRead);
if(nSize != nBytesRead)
assert(false);
+ length = nBytesRead;
+
pBinary[i] = binary;
}
@@ -467,7 +469,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
fprintf(stderr, "Create kernel from binary\n");
gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary( gpuInfo->mpContext,numDevices,
- mpArryDevsID.get(), &length, (const unsigned char**) &pBinary,
+ mpArryDevsID.get(), &length, (const unsigned char**) pBinary.get(),
&binary_status, &clStatus );
if(clStatus != CL_SUCCESS)
{
commit 81500c22cae0dbd79bb9fbde077ccad752109db1
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Tue Sep 17 19:31:41 2013 +0200
use $UserInstallation/cache to cache opencl kernels
Change-Id: I2ba1fb6172cfd0c725a45d4506b46e8f04a33093
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 45dedc2..1a19451 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -18,6 +18,7 @@
#include <boost/scoped_array.hpp>
#include "sal/config.h"
+#include <osl/file.hxx>
#include "oclkernels.hxx"
#include <stdio.h>
commit 801348462cbca07fd92f63c24eb8b92109a3e0e4
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Tue Sep 17 23:04:54 2013 +0200
read the binary files from the cached folder
Change-Id: Ifc5084a3d8ef661e3b3fca7fbc76e1c0c62b2056
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 1ca38fd..45dedc2 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -18,7 +18,6 @@
#include <boost/scoped_array.hpp>
#include "sal/config.h"
-#include <osl/file.hxx>
#include "oclkernels.hxx"
#include <stdio.h>
@@ -216,45 +215,42 @@ OString createFileName(cl_device_id deviceId, const char* clFileName)
}
-int OpenclDevice::binaryGenerated( const char * clFileName, FILE ** fhandle )
+std::vector<boost::shared_ptr<osl::File> > OpenclDevice::binaryGenerated( const char * clFileName, cl_context context )
{
- unsigned int i = 0;
- cl_int clStatus;
- int status = 0;
- FILE *fd = NULL;
cl_uint numDevices=0;
- 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++ )
+ std::vector<boost::shared_ptr<osl::File> > aGeneratedFiles;
+ cl_int clStatus = clGetContextInfo( context, CL_CONTEXT_NUM_DEVICES,
+ sizeof(numDevices), &numDevices, NULL );
+
+ if(clStatus != CL_SUCCESS)
+ return aGeneratedFiles;
+
+
+ // grab the handles to all of the devices in the context.
+ boost::scoped_array<cl_device_id> mpArryDevsID(new cl_device_id[numDevices]);
+ clStatus = clGetContextInfo( context, CL_CONTEXT_DEVICES,
+ sizeof( cl_device_id ) * numDevices, mpArryDevsID.get(), NULL );
+
+ if(clStatus != CL_SUCCESS)
+ return aGeneratedFiles;
+
+ for ( size_t i = 0; i < numDevices; i++ )
{
- if ( gpuEnv.mpArryDevsID[i] != 0 )
+ if ( mpArryDevsID[i] != 0 )
{
OString fileName = createFileName(gpuEnv.mpArryDevsID[i], clFileName);
- fd = fopen( fileName.getStr(), "rb" );
- status = ( fd != NULL ) ? 1 : 0;
+ osl::File* pNewFile = new osl::File(rtl::OStringToOUString(fileName, RTL_TEXTENCODING_UTF8));
+ if(pNewFile->open(osl_File_OpenFlag_Read) == osl::FileBase::E_None)
+ aGeneratedFiles.push_back(boost::shared_ptr<osl::File>(pNewFile));
+ else
+ {
+ delete pNewFile;
+ break;
+ }
}
}
- if ( fd != NULL )
- {
- *fhandle = fd;
- }
- return status;
+ return aGeneratedFiles;
}
@@ -418,10 +414,8 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
{
cl_int clStatus = 0;
size_t length;
- char *binary;
- int b_error, binary_status, binaryExisted, idx;
+ int binary_status, idx;
cl_uint numDevices;
- FILE *fd;
const char* filename = "kernel.cl";
fprintf(stderr, "compileKernelFile ... \n");
if ( cachedOfKernerPrg(gpuInfo, filename) == 1 )
@@ -431,49 +425,57 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
idx = gpuInfo->mnFileCount;
- binaryExisted = 0;
- if ( ( binaryExisted = binaryGenerated( filename, &fd ) ) == 1 )
- {
- clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES,
- sizeof(numDevices), &numDevices, NULL );
- CHECK_OPENCL( clStatus, "clGetContextInfo" );
+ clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES,
+ sizeof(numDevices), &numDevices, NULL );
+ CHECK_OPENCL( clStatus, "clGetContextInfo" );
- boost::scoped_array<cl_device_id> mpArryDevsID(new cl_device_id[numDevices]);
+ std::vector<boost::shared_ptr<osl::File> > aGeneratedFiles = binaryGenerated(
+ filename, gpuInfo->mpContext );
- 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 )
+ bool bBinaryExisted = false;
+ if (aGeneratedFiles.size() == numDevices)
+ {
+ bBinaryExisted = true;
+ boost::scoped_array<char*> pBinary(new char*[numDevices]);
+ for(size_t i = 0; i < numDevices; ++i)
{
- return 0;
+ sal_uInt64 nSize;
+ aGeneratedFiles[i]->getSize(nSize);
+ char* binary = new char[nSize];
+ sal_uInt64 nBytesRead;
+ aGeneratedFiles[i]->read(binary, nSize, nBytesRead);
+ if(nSize != nBytesRead)
+ assert(false);
+
+ pBinary[i] = binary;
}
- memset( binary, 0, length + 2 );
- b_error |= fread( binary, 1, length, fd ) != length;
-
-
- fclose( fd );
- fd = NULL;
// grab the handles to all of the devices in the context.
+ boost::scoped_array<cl_device_id> mpArryDevsID(new cl_device_id[numDevices]);
clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_DEVICES,
sizeof( cl_device_id ) * numDevices, mpArryDevsID.get(), NULL );
+
+ if(clStatus != CL_SUCCESS)
+ {
+ for(size_t i = 0; i < numDevices; ++i)
+ {
+ delete[] pBinary[i];
+ }
+ }
CHECK_OPENCL( clStatus, "clGetContextInfo" );
fprintf(stderr, "Create kernel from binary\n");
gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary( gpuInfo->mpContext,numDevices,
- mpArryDevsID.get(), &length, (const unsigned char**) &binary,
+ mpArryDevsID.get(), &length, (const unsigned char**) &pBinary,
&binary_status, &clStatus );
+ if(clStatus != CL_SUCCESS)
+ {
+ for(size_t i = 0; i < numDevices; ++i)
+ {
+ delete[] pBinary[i];
+ }
+ }
CHECK_OPENCL( clStatus, "clCreateProgramWithBinary" );
-
- free( binary );
}
else
{
@@ -558,7 +560,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
strcpy( gpuEnv.mArryKnelSrcFile[idx], filename );
- if ( binaryExisted == 0 )
+ if ( !bBinaryExisted )
generatBinFromKernelSource( gpuEnv.mpArryPrograms[idx], filename );
gpuInfo->mnFileCount += 1;
diff --git a/sc/source/core/opencl/openclwrapper.hxx b/sc/source/core/opencl/openclwrapper.hxx
index 7bcf5c4..c81c313 100644
--- a/sc/source/core/opencl/openclwrapper.hxx
+++ b/sc/source/core/opencl/openclwrapper.hxx
@@ -13,6 +13,9 @@
#include <config_features.h>
#include <formula/opcode.hxx>
#include <sal/detail/log.h>
+#include <osl/file.hxx>
+#include <vector>
+#include <boost/shared_ptr.hpp>
#include <cassert>
#include "platforminfo.hxx"
@@ -174,7 +177,7 @@ public:
static int cachedOfKernerPrg( const GPUEnv *gpuEnvCached, const char * clFileName );
static int generatBinFromKernelSource( cl_program program, const char * clFileName );
static int writeBinaryToFile( const OString& rName, const char* birary, size_t numBytes );
- static int binaryGenerated( const char * clFileName, FILE ** fhandle );
+ static std::vector<boost::shared_ptr<osl::File> > binaryGenerated( const char * clFileName, cl_context context);
static int compileKernelFile( const char *filename, GPUEnv *gpuInfo, const char *buildOption );
static int initOpenclAttr( OpenCLEnv * env );
commit 1f71e29cdece933258d67da172e4cea4d2fb27c4
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Tue Sep 17 20:30:53 2013 +0200
replace one more file write with location in profile directory
Change-Id: I1774f66d018923c6f892d21acf8a81330628cd98
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 98ef0d2..1ca38fd 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -419,10 +419,9 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
cl_int clStatus = 0;
size_t length;
char *binary;
- const char *source;
int b_error, binary_status, binaryExisted, idx;
cl_uint numDevices;
- FILE *fd, *fd1;
+ FILE *fd;
const char* filename = "kernel.cl";
fprintf(stderr, "compileKernelFile ... \n");
if ( cachedOfKernerPrg(gpuInfo, filename) == 1 )
@@ -432,8 +431,6 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
idx = gpuInfo->mnFileCount;
- source = kernel_src;
-
binaryExisted = 0;
if ( ( binaryExisted = binaryGenerated( filename, &fd ) ) == 1 )
{
@@ -483,8 +480,9 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
// create a CL program using the kernel source
fprintf(stderr, "Create kernel from source\n");
size_t source_size[1];
- source_size[0] = strlen( source );
- gpuEnv.mpArryPrograms[idx] = clCreateProgramWithSource( gpuEnv.mpContext, 1, &source,
+
+ source_size[0] = strlen( kernel_src );
+ gpuEnv.mpArryPrograms[idx] = clCreateProgramWithSource( gpuEnv.mpContext, 1, &kernel_src,
source_size, &clStatus);
CHECK_OPENCL( clStatus, "clCreateProgramWithSource" );
}
@@ -544,12 +542,16 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
return 0;
}
- fd1 = fopen( "kernel-build.log", "w+" );
- if ( fd1 != NULL )
- {
- fwrite( buildLog.get(), sizeof(char), length, fd1 );
- fclose( fd1 );
- }
+ OString aBuildLogFileURL = maCacheFolder + "kernel-build.log";
+ osl::File aBuildLogFile(rtl::OStringToOUString(aBuildLogFileURL, RTL_TEXTENCODING_UTF8));
+ osl::FileBase::RC status = aBuildLogFile.open(
+ osl_File_OpenFlag_Write | osl_File_OpenFlag_Create );
+
+ if(status != osl::FileBase::E_None)
+ return 0;
+
+ sal_uInt64 nBytesWritten = 0;
+ aBuildLogFile.write( buildLog.get(), length, nBytesWritten );
return 0;
}
commit 742aa94f3399dbcaa9f6067bb34774f99733f847
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Tue Sep 17 19:49:20 2013 +0200
fix another memory leak
Change-Id: I6761bcc137934b02815ce10d43f3bc9bee7a1b90
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 2ec2046..98ef0d2 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -418,7 +418,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
{
cl_int clStatus = 0;
size_t length;
- char *buildLog = NULL, *binary;
+ char *binary;
const char *source;
int b_error, binary_status, binaryExisted, idx;
cl_uint numDevices;
@@ -526,20 +526,17 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
printf("opencl create build log fail\n");
return 0;
}
- buildLog = (char*) malloc( length );
- if ( buildLog == (char*) NULL )
- {
- return 0;
- }
+
+ boost::scoped_array<char> buildLog(new char[length]);
if ( !gpuInfo->mnIsUserCreated )
{
clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpArryDevsID[0],
- CL_PROGRAM_BUILD_LOG, length, buildLog, &length );
+ CL_PROGRAM_BUILD_LOG, length, buildLog.get(), &length );
}
else
{
clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpDevID,
- CL_PROGRAM_BUILD_LOG, length, buildLog, &length );
+ CL_PROGRAM_BUILD_LOG, length, buildLog.get(), &length );
}
if ( clStatus != CL_SUCCESS )
{
@@ -550,11 +547,10 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
fd1 = fopen( "kernel-build.log", "w+" );
if ( fd1 != NULL )
{
- fwrite( buildLog, sizeof(char), length, fd1 );
+ fwrite( buildLog.get(), sizeof(char), length, fd1 );
fclose( fd1 );
}
- free( buildLog );
return 0;
}
commit 6b621cffa726f627750eb54093029cad98e497b6
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Tue Sep 17 19:36:09 2013 +0200
fix another memory leak
Change-Id: I31359f121fa858dfc1868d74a2f827e5332592c3
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 3610659..2ec2046 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -420,10 +420,8 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
size_t length;
char *buildLog = NULL, *binary;
const char *source;
- size_t source_size[1];
int b_error, binary_status, binaryExisted, idx;
cl_uint numDevices;
- cl_device_id *mpArryDevsID;
FILE *fd, *fd1;
const char* filename = "kernel.cl";
fprintf(stderr, "compileKernelFile ... \n");
@@ -436,7 +434,6 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
source = kernel_src;
- source_size[0] = strlen( source );
binaryExisted = 0;
if ( ( binaryExisted = binaryGenerated( filename, &fd ) ) == 1 )
{
@@ -444,11 +441,7 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
sizeof(numDevices), &numDevices, NULL );
CHECK_OPENCL( clStatus, "clGetContextInfo" );
- mpArryDevsID = (cl_device_id*) malloc( sizeof(cl_device_id) * numDevices );
- if ( mpArryDevsID == NULL )
- {
- return 0;
- }
+ boost::scoped_array<cl_device_id> mpArryDevsID(new cl_device_id[numDevices]);
b_error = 0;
length = 0;
@@ -474,23 +467,23 @@ int OpenclDevice::compileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
fd = NULL;
// 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 );
+ sizeof( cl_device_id ) * numDevices, mpArryDevsID.get(), NULL );
CHECK_OPENCL( clStatus, "clGetContextInfo" );
fprintf(stderr, "Create kernel from binary\n");
gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary( gpuInfo->mpContext,numDevices,
- mpArryDevsID, &length, (const unsigned char**) &binary,
+ mpArryDevsID.get(), &length, (const unsigned char**) &binary,
&binary_status, &clStatus );
CHECK_OPENCL( clStatus, "clCreateProgramWithBinary" );
free( binary );
- free( mpArryDevsID );
- mpArryDevsID = NULL;
}
else
{
// create a CL program using the kernel source
fprintf(stderr, "Create kernel from source\n");
+ size_t source_size[1];
+ source_size[0] = strlen( source );
gpuEnv.mpArryPrograms[idx] = clCreateProgramWithSource( gpuEnv.mpContext, 1, &source,
source_size, &clStatus);
CHECK_OPENCL( clStatus, "clCreateProgramWithSource" );
commit 25b9ca83d2be00ca0c4260fb4c47884c13c95aab
Author: Markus Mohrhard <markus.mohrhard at googlemail.com>
Date: Tue Sep 17 19:31:41 2013 +0200
use $UserInstallation/cache to cache opencl kernels
Change-Id: I2ba1fb6172cfd0c725a45d4506b46e8f04a33093
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 684abe6..3610659 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -7,14 +7,18 @@
* file, You can obtain one at http://mozilla.org/MPL/2.0/.
*/
+#include <config_folders.h>
+
#include "openclwrapper.hxx"
#include <rtl/ustring.hxx>
#include <rtl/strbuf.hxx>
#include <rtl/digest.h>
+#include <rtl/bootstrap.hxx>
#include <boost/scoped_array.hpp>
#include "sal/config.h"
+#include <osl/file.hxx>
#include "oclkernels.hxx"
#include <stdio.h>
@@ -76,9 +80,20 @@ OString generateHashForSource()
return aBuffer.makeStringAndClear();
}
+OString getCacheFolder()
+{
+ OUString url("${$BRAND_BASE_DIR/" LIBO_ETC_FOLDER "/" SAL_CONFIGFILE("bootstrap") ":UserInstallation}/cache/");
+ rtl::Bootstrap::expandMacros(url);
+
+ osl::Directory::create(url);
+
+ return rtl::OUStringToOString(url, RTL_TEXTENCODING_UTF8);
+}
+
}
OString OpenclDevice::maSourceHash = generateHashForSource();
+OString OpenclDevice::maCacheFolder = getCacheFolder();
int OpenclDevice::releaseOpenclRunEnv()
{
@@ -195,7 +210,8 @@ OString createFileName(cl_device_id deviceId, const char* clFileName)
char deviceName[DEVICE_NAME_LENGTH] = {0};
clGetDeviceInfo(deviceId, CL_DEVICE_NAME,
sizeof(deviceName), deviceName, NULL);
- return fileName + "-" + deviceName + "-" + OpenclDevice::maSourceHash + ".bin";
+ return OpenclDevice::maCacheFolder + fileName + "-" +
+ deviceName + "-" + OpenclDevice::maSourceHash + ".bin";
}
}
@@ -242,17 +258,19 @@ int OpenclDevice::binaryGenerated( const char * clFileName, FILE ** fhandle )
}
-int OpenclDevice::writeBinaryToFile( const OString& rFileName, const char* birary, size_t numBytes )
+int OpenclDevice::writeBinaryToFile( const OString& rFileName, const char* binary, size_t numBytes )
{
- FILE *output = NULL;
- output = fopen( rFileName.getStr(), "wb" );
- if ( output == NULL )
- {
+ osl::File file(rtl::OStringToOUString(rFileName, RTL_TEXTENCODING_UTF8));
+ osl::FileBase::RC status = file.open(
+ osl_File_OpenFlag_Write | osl_File_OpenFlag_Create );
+
+ if(status != osl::FileBase::E_None)
return 0;
- }
- fwrite( birary, sizeof(char), numBytes, output );
- fclose( output );
+ sal_uInt64 nBytesWritten = 0;
+ file.write( binary, numBytes, nBytesWritten );
+
+ assert(numBytes == nBytesWritten);
return 1;
diff --git a/sc/source/core/opencl/openclwrapper.hxx b/sc/source/core/opencl/openclwrapper.hxx
index 5a89e6b..7bcf5c4 100644
--- a/sc/source/core/opencl/openclwrapper.hxx
+++ b/sc/source/core/opencl/openclwrapper.hxx
@@ -163,6 +163,8 @@ public:
static GPUEnv gpuEnv;
static int isInited;
static OString maSourceHash;
+ static OString maCacheFolder;
+
static int registOpenclKernel();
static int releaseOpenclRunEnv();
static int initOpenclRunEnv( GPUEnv *gpu );
commit f708ed03f4f9fa203080b11e03ad0c9ad8e2adfa
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 23:34:11 2013 -0400
In xlsx, a better heuristic is based on worksheet dimension.
And start OpenCL kernel compilation right away if the dimension is
large enough to make the sheet import slow.
Change-Id: Iee3e8dc1f0cb64eb1d8778db19c6edd73e971136
diff --git a/sc/source/filter/oox/formulabuffer.cxx b/sc/source/filter/oox/formulabuffer.cxx
index cc8a50a..21e383f 100644
--- a/sc/source/filter/oox/formulabuffer.cxx
+++ b/sc/source/filter/oox/formulabuffer.cxx
@@ -65,7 +65,6 @@ void FormulaBuffer::finalizeImport()
FormulaDataMap::iterator cellIt = maCellFormulas.find( nTab );
if ( cellIt != maCellFormulas.end() )
{
- compileOpenCLKernels();
applyCellFormulas( cellIt->second );
}
diff --git a/sc/source/filter/oox/worksheetfragment.cxx b/sc/source/filter/oox/worksheetfragment.cxx
index b8ad658..93cacd7 100644
--- a/sc/source/filter/oox/worksheetfragment.cxx
+++ b/sc/source/filter/oox/worksheetfragment.cxx
@@ -511,7 +511,16 @@ void WorksheetFragment::importDimension( const AttributeList& rAttribs )
be set. If the cell A1 exists, the used area will be updated while
importing the cell. */
if( (aRange.EndColumn > 0) || (aRange.EndRow > 0) )
+ {
extendUsedArea( aRange );
+
+ // Start pre-compiling OpenCL kernels if we have a large number of
+ // cells to import. The current threshold is 100,000.
+ double fCellCount = aRange.EndColumn - aRange.StartColumn + 1.0;
+ fCellCount *= aRange.EndRow - aRange.StartRow + 1.0;
+ if (fCellCount > 100000.0)
+ compileOpenCLKernels();
+ }
}
void WorksheetFragment::importSheetFormatPr( const AttributeList& rAttribs )
commit aaf8ec3a05a71e977287028643da6f8ca0a0548a
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 20:57:48 2013 -0400
Be explicit on what kind of "kernels" we are compiling.
Change-Id: I2415a192d98b94e86c08c138ad39cfdbf36ea283
diff --git a/sc/inc/formulagroup.hxx b/sc/inc/formulagroup.hxx
index c6f32ac..e7c13cd 100644
--- a/sc/inc/formulagroup.hxx
+++ b/sc/inc/formulagroup.hxx
@@ -56,7 +56,7 @@ class SC_DLLPUBLIC FormulaGroupInterpreter
static FormulaGroupInterpreter *getStatic();
static void fillOpenCLInfo(std::vector<OpenclPlatformInfo>& rPlatforms);
static void switchOpenCLDevice(const OUString& rDeviceId, bool bAutoSelect);
- static void compileKernels();
+ static void compileOpenCLKernels();
virtual ScMatrixRef inverseMatrix(const ScMatrix& rMat) = 0;
virtual bool interpret(ScDocument& rDoc, const ScAddress& rTopPos, const ScFormulaCellGroupRef& xGroup, ScTokenArray& rCode) = 0;
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index a311179..ab7076d 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -1095,9 +1095,9 @@ SAL_DLLPUBLIC_EXPORT bool SAL_CALL switchOpenClDevice(const OUString* pDeviceId,
return sc::opencl::switchOpenclDevice(pDeviceId, bAutoSelect);
}
-SAL_DLLPUBLIC_EXPORT void compileKernels(const OUString* pDeviceId)
+SAL_DLLPUBLIC_EXPORT void compileOpenCLKernels(const OUString* pDeviceId)
{
- sc::opencl::compileKernels(pDeviceId);
+ sc::opencl::compileOpenCLKernels(pDeviceId);
}
}
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 78d3bef..684abe6 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -2744,7 +2744,7 @@ bool switchOpenclDevice(const OUString* pDevice, bool bAutoSelect)
return !OpenclDevice::initOpenclRunEnv(0);
}
-void compileKernels(const OUString* pDeviceId)
+void compileOpenCLKernels(const OUString* pDeviceId)
{
if (!pDeviceId)
return;
diff --git a/sc/source/core/opencl/openclwrapper.hxx b/sc/source/core/opencl/openclwrapper.hxx
index dfa8fbb..5a89e6b 100644
--- a/sc/source/core/opencl/openclwrapper.hxx
+++ b/sc/source/core/opencl/openclwrapper.hxx
@@ -266,7 +266,7 @@ const std::vector<OpenclPlatformInfo>& fillOpenCLInfo();
*/
bool switchOpenclDevice(const OUString* pDeviceId, bool bAutoSelect);
-void compileKernels(const OUString* pDeviceId);
+void compileOpenCLKernels(const OUString* pDeviceId);
}}
diff --git a/sc/source/core/tool/formulagroup.cxx b/sc/source/core/tool/formulagroup.cxx
index 6a20a0c..2f85ee5 100644
--- a/sc/source/core/tool/formulagroup.cxx
+++ b/sc/source/core/tool/formulagroup.cxx
@@ -338,7 +338,7 @@ typedef FormulaGroupInterpreter* (*__createFormulaGroupOpenCLInterpreter)(void);
typedef size_t (*__getOpenCLPlatformCount)(void);
typedef void (*__fillOpenCLInfo)(OpenclPlatformInfo*, size_t);
typedef bool (*__switchOpenClDevice)(const OUString*, bool);
-typedef void (*__compileKernels)(const OUString*);
+typedef void (*__compileOpenCLKernels)(const OUString*);
#endif
@@ -458,7 +458,7 @@ void FormulaGroupInterpreter::switchOpenCLDevice(const OUString& rDeviceId, bool
#endif
}
-void FormulaGroupInterpreter::compileKernels()
+void FormulaGroupInterpreter::compileOpenCLKernels()
{
const ScCalcConfig& rConfig = ScInterpreter::GetGlobalConfig();
if (!rConfig.mbOpenCLEnabled)
@@ -469,11 +469,11 @@ void FormulaGroupInterpreter::compileKernels()
if (!pModule)
return;
- oslGenericFunction fn = pModule->getFunctionSymbol("compileKernels");
+ oslGenericFunction fn = pModule->getFunctionSymbol("compileOpenCLKernels");
if (!fn)
return;
- reinterpret_cast<__compileKernels>(fn)(&rConfig.maOpenCLDevice);
+ reinterpret_cast<__compileOpenCLKernels>(fn)(&rConfig.maOpenCLDevice);
}
void FormulaGroupInterpreter::generateRPNCode(ScDocument& rDoc, const ScAddress& rPos, ScTokenArray& rCode)
diff --git a/sc/source/filter/ftools/clkernelthread.cxx b/sc/source/filter/ftools/clkernelthread.cxx
index f6d8c63..dcb1e74 100644
--- a/sc/source/filter/ftools/clkernelthread.cxx
+++ b/sc/source/filter/ftools/clkernelthread.cxx
@@ -20,7 +20,7 @@ CLBuildKernelThread::~CLBuildKernelThread() {}
void CLBuildKernelThread::execute()
{
- sc::FormulaGroupInterpreter::compileKernels();
+ sc::FormulaGroupInterpreter::compileOpenCLKernels();
}
}
commit 0dc679a4285f466c0ec7650b2e31e242ea466988
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 20:53:24 2013 -0400
Add a heuristic to avoid always pre-compile OpenCL kernels (for xls).
If we still have 3 MB worth of stream to parse after the first formula
record, we pre-compile OpenCL kernels in a separate thread.
Change-Id: I129d95c26f26e4b1eb989bc2a1764a3f87f71fcb
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 2b2f8f7..78d3bef 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -2770,7 +2770,6 @@ void compileKernels(const OUString* pDeviceId)
if (nStatus != CL_SUCCESS)
r.mpKernel = NULL;
}
-
}
}}
diff --git a/sc/source/filter/excel/excform.cxx b/sc/source/filter/excel/excform.cxx
index 49ca6ca..b557de5 100644
--- a/sc/source/filter/excel/excform.cxx
+++ b/sc/source/filter/excel/excform.cxx
@@ -136,11 +136,19 @@ void ImportExcel::Formula(
}
}
- if (!mxCLKernelThread.is())
+ if (mbRunCLKernelThread && !mxCLKernelThread.is())
{
- mxCLKernelThread.set(new sc::CLBuildKernelThread);
- mxCLKernelThread->launch();
+ double fRemainSize = maStrm.GetSvStreamSize() - maStrm.GetSvStreamPos();
+ fRemainSize /= 1048576; // Switch to MB.
+ if (fRemainSize > 3.0)
+ {
+ // We still have more than 3MB of stream to parse. Start a thread to pre-compile OpenCL kernels.
+ mxCLKernelThread.set(new sc::CLBuildKernelThread);
+ mxCLKernelThread->launch();
+ }
+ mbRunCLKernelThread = false;
}
+
ConvErr eErr = pFormConv->Convert( pResult, maStrm, nFormLen, true, FT_CellFormula);
ScFormulaCell* pCell = NULL;
diff --git a/sc/source/filter/excel/impop.cxx b/sc/source/filter/excel/impop.cxx
index e03980b..16c4ac5 100644
--- a/sc/source/filter/excel/impop.cxx
+++ b/sc/source/filter/excel/impop.cxx
@@ -116,8 +116,9 @@ ImportExcel::ImportExcel( XclImpRootData& rImpData, SvStream& rStrm ):
maScOleSize( ScAddress::INITIALIZE_INVALID ),
mnLastRefIdx( 0 ),
mnIxfeIndex( 0 ),
- mbBiff2HasXfs( false ),
- mbBiff2HasXfsValid( false )
+ mbBiff2HasXfs(false),
+ mbBiff2HasXfsValid(false),
+ mbRunCLKernelThread(true)
{
nBdshtTab = 0;
diff --git a/sc/source/filter/inc/imp_op.hxx b/sc/source/filter/inc/imp_op.hxx
index 613ae30..6e44f31 100644
--- a/sc/source/filter/inc/imp_op.hxx
+++ b/sc/source/filter/inc/imp_op.hxx
@@ -106,8 +106,6 @@ protected:
sal_Int16 mnLastRefIdx;
sal_uInt16 mnIxfeIndex; /// Current XF identifier from IXFE record.
- bool mbBiff2HasXfs; /// Select XF formatting or direct formatting in BIFF2.
- bool mbBiff2HasXfsValid; /// False = mbBiff2HasXfs is undetermined yet.
SCTAB nBdshtTab; // Counter fuer Boundsheet
ScFormulaCell* pLastFormCell; // fuer String-Records
@@ -115,6 +113,10 @@ protected:
sal_Bool bTabTruncated; // wenn Bereichsueberschreitung zum
// Abschneiden von Zellen fuehrt
+ bool mbBiff2HasXfs:1; /// Select XF formatting or direct formatting in BIFF2.
+ bool mbBiff2HasXfsValid:1; /// False = mbBiff2HasXfs is undetermined yet.
+ bool mbRunCLKernelThread:1;
+
// Record-Funktionen
void ReadFileSharing();
commit a00dbd951aff752b40ec7892cdf69f125e7a0137
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 17:26:45 2013 -0400
Oops this wasn't meant to be checked in.
Change-Id: I6eccbd38a44dfaec66718df210884067beaaf30e
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 7bcd390..2b2f8f7 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -43,49 +43,6 @@
#define DEVICE_NAME_LENGTH 1024
-
-#include <stdio.h>
-#include <string>
-#include <sys/time.h>
-
-namespace {
-
-class stack_printer
-{
-public:
- explicit stack_printer(const char* msg) :
- msMsg(msg)
- {
- fprintf(stdout, "%s: --begin\n", msMsg.c_str());
- mfStartTime = getTime();
- }
-
- ~stack_printer()
- {
- double fEndTime = getTime();
- fprintf(stdout, "%s: --end (duration: %g sec)\n", msMsg.c_str(), (fEndTime - mfStartTime));
- }
-
- void printTime(int line) const
- {
- double fEndTime = getTime();
- fprintf(stdout, "%s: --(%d) (duration: %g sec)\n", msMsg.c_str(), line, (fEndTime - mfStartTime));
- }
-
-private:
- double getTime() const
- {
- timeval tv;
- gettimeofday(&tv, NULL);
- return tv.tv_sec + tv.tv_usec / 1000000.0;
- }
-
- ::std::string msMsg;
- double mfStartTime;
-};
-
-}
-
using namespace std;
namespace sc { namespace opencl {
@@ -2795,8 +2752,6 @@ void compileKernels(const OUString* pDeviceId)
if (pDeviceId->isEmpty())
return;
- stack_printer __stack_printer__("sc/opencl::compileKernels");
- fprintf(stdout, "opencl::compileKernels: device = '%s'\n", rtl::OUStringToOString(*pDeviceId, RTL_TEXTENCODING_UTF8).getStr());
if (!switchOpenclDevice(pDeviceId, false))
return;
commit 99837d78225d994e810d69abb6f65058a81c0d33
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 17:22:38 2013 -0400
Fix illegal memory access between malloc and memset.
Change-Id: I6f040e48a1a60b8f4fc1f2424a24ff362de9a31e
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index 139a219..a311179 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -455,66 +455,54 @@ double * agency::calculate( int nOclOp,int rowSize,OclCalc &ocl_calc,uint *npOcl
{
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 );
+ rResult = (double *)malloc( sizeof(double) * rowSize );
memset(rResult,0,rowSize);
ocl_calc.oclHostArithmeticStash64Bits( "oclSignedAdd",dpLeftData,dpRightData,rResult,temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ formulaInterprt->srdDataPush( new SourceData( rResult,rowSize ) );
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 );
+ rResult = ( double * )malloc( sizeof(double) * rowSize );
memset( rResult,0,rowSize );
ocl_calc.oclHostArithmeticStash64Bits( "oclSignedSub",dpLeftData,dpRightData,rResult,temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData(rResult,nDataSize) );
+ formulaInterprt->srdDataPush( new SourceData(rResult,rowSize) );
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 );
+ rResult = (double *)malloc( sizeof(double) * rowSize );
memset( rResult,0,rowSize );
ocl_calc.oclHostArithmeticStash64Bits( "oclSignedMul",dpLeftData,dpRightData,rResult,temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ formulaInterprt->srdDataPush( new SourceData( rResult,rowSize ) );
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 );
+ rResult = ( double * )malloc( sizeof(double) * rowSize );
memset( rResult,0,rowSize );
ocl_calc.oclHostArithmeticStash64Bits( "oclSignedDiv",dpLeftData,dpRightData,rResult,temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ formulaInterprt->srdDataPush( new SourceData( rResult,rowSize ) );
break;
}
case ocMax:
@@ -567,66 +555,54 @@ double * agency::calculate( int nOclOp,int rowSize,OclCalc &ocl_calc,uint *npOcl
{
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 );
+ rResult = (double *)malloc( sizeof(double) * rowSize );
memset(rResult,0,rowSize);
ocl_calc.oclHostArithmeticStash32Bits( "oclSignedAdd", dpLeftData, dpRightData, rResult, temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData(rResult, nDataSize) );
+ formulaInterprt->srdDataPush( new SourceData(rResult, rowSize) );
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 );
+ rResult = (double *)malloc( sizeof(double) * rowSize );
memset( rResult, 0, rowSize );
ocl_calc.oclHostArithmeticStash32Bits( "oclSignedSub", dpLeftData, dpRightData, rResult, temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData( rResult,nDataSize ) );
+ formulaInterprt->srdDataPush( new SourceData( rResult, rowSize ) );
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 );
+ rResult = (double *)malloc(sizeof(double) * rowSize );
memset( rResult, 0, rowSize );
ocl_calc.oclHostArithmeticStash32Bits( "oclSignedMul", dpLeftData, dpRightData, rResult, temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData( rResult, nDataSize ) );
+ formulaInterprt->srdDataPush( new SourceData( rResult, rowSize ) );
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 );
+ rResult = (double *)malloc( sizeof(double) * rowSize );
memset( rResult, 0, rowSize );
ocl_calc.oclHostArithmeticStash32Bits( "oclSignedDiv", dpLeftData, dpRightData, rResult, temp->getDataSize() );
- formulaInterprt->srdDataPush( new SourceData(rResult, nDataSize) );
+ formulaInterprt->srdDataPush( new SourceData(rResult, rowSize) );
break;
}
case ocMax:
@@ -636,7 +612,7 @@ double * agency::calculate( int nOclOp,int rowSize,OclCalc &ocl_calc,uint *npOcl
nDataSize = temp->getDataSize();
dpOclSrcData = temp->getDouleData();
double *rResult = NULL; // Point to the output data from GPU
- rResult = (double *)malloc(sizeof(double) * nDataSize );
+ rResult = (double *)malloc(sizeof(double) * rowSize );
memset(rResult,0,rowSize);
ocl_calc.oclHostFormulaStash32Bits( "oclFormulaMax", dpOclSrcData, npOclStartPos, npOclEndPos, rResult,nDataSize, rowSize );
formulaInterprt->srdDataPush( new SourceData( rResult, rowSize ) );
@@ -649,7 +625,7 @@ double * agency::calculate( int nOclOp,int rowSize,OclCalc &ocl_calc,uint *npOcl
nDataSize = temp->getDataSize();
dpOclSrcData = temp->getDouleData();
double *rResult = NULL; // Point to the output data from GPU
- rResult = (double *)malloc( sizeof(double) * nDataSize );
+ rResult = (double *)malloc( sizeof(double) * rowSize );
memset( rResult, 0, rowSize );
ocl_calc.oclHostFormulaStash32Bits( "oclFormulaMin", dpOclSrcData, npOclStartPos, npOclEndPos, rResult, nDataSize, rowSize );
formulaInterprt->srdDataPush( new SourceData( rResult, rowSize) );
@@ -662,7 +638,7 @@ double * agency::calculate( int nOclOp,int rowSize,OclCalc &ocl_calc,uint *npOcl
nDataSize = temp->getDataSize();
dpOclSrcData = temp->getDouleData();
double *rResult = NULL; // Point to the output data from GPU
- rResult = (double *)malloc( sizeof(double) * nDataSize );
+ rResult = (double *)malloc( sizeof(double) * rowSize );
memset( rResult, 0, rowSize);
ocl_calc.oclHostFormulaStash32Bits( "oclFormulaAverage", dpOclSrcData, npOclStartPos, npOclEndPos, rResult, nDataSize, rowSize );
formulaInterprt->srdDataPush( new SourceData( rResult, rowSize) );
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 2b2f8f7..7bcd390 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -43,6 +43,49 @@
#define DEVICE_NAME_LENGTH 1024
+
+#include <stdio.h>
+#include <string>
+#include <sys/time.h>
+
+namespace {
+
+class stack_printer
+{
+public:
+ explicit stack_printer(const char* msg) :
+ msMsg(msg)
+ {
+ fprintf(stdout, "%s: --begin\n", msMsg.c_str());
+ mfStartTime = getTime();
+ }
+
+ ~stack_printer()
+ {
+ double fEndTime = getTime();
+ fprintf(stdout, "%s: --end (duration: %g sec)\n", msMsg.c_str(), (fEndTime - mfStartTime));
+ }
+
+ void printTime(int line) const
+ {
+ double fEndTime = getTime();
+ fprintf(stdout, "%s: --(%d) (duration: %g sec)\n", msMsg.c_str(), line, (fEndTime - mfStartTime));
+ }
+
+private:
+ double getTime() const
+ {
+ timeval tv;
+ gettimeofday(&tv, NULL);
+ return tv.tv_sec + tv.tv_usec / 1000000.0;
+ }
+
+ ::std::string msMsg;
+ double mfStartTime;
+};
+
+}
+
using namespace std;
namespace sc { namespace opencl {
@@ -2752,6 +2795,8 @@ void compileKernels(const OUString* pDeviceId)
if (pDeviceId->isEmpty())
return;
+ stack_printer __stack_printer__("sc/opencl::compileKernels");
+ fprintf(stdout, "opencl::compileKernels: device = '%s'\n", rtl::OUStringToOString(*pDeviceId, RTL_TEXTENCODING_UTF8).getStr());
if (!switchOpenclDevice(pDeviceId, false))
return;
commit 491b9c848e7b5c46f412d6e53c21d370a1cfe989
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 16:48:20 2013 -0400
Bail out early on error.
Change-Id: I1844284ff829f6af6c2cfe10935cfc1eef155deb
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index 5697b1b..139a219 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -308,7 +308,14 @@ bool FormulaGroupInterpreterOpenCL::chooseFunction( OclCalc &ocl_calc, double *&
pGroundWaterDataArray=NULL;
delta = mdpSvdouble[0];
}
- ocl_calc.oclGroundWaterGroup( mnOperatorGroup,mnOperatorCount,pGroundWaterDataArray,pArrayToSubtractOneElementFrom,nSrcData,mnRowSize,delta,mnpOclStartPos,mnpOclEndPos,dpResult);
+
+ bool bSuccess = ocl_calc.oclGroundWaterGroup(
+ mnOperatorGroup, mnOperatorCount, pGroundWaterDataArray,
+ pArrayToSubtractOneElementFrom, nSrcData, mnRowSize, delta,
+ mnpOclStartPos, mnpOclEndPos, dpResult);
+
+ if (!bSuccess)
+ return false;
}
else if( isStockHistory() )
{
@@ -332,14 +339,19 @@ bool FormulaGroupInterpreterOpenCL::chooseFunction( OclCalc &ocl_calc, double *&
{
if (!ocl_calc.createFormulaBuf64Bits(nSrcDataSize, mnRowSize))
return false;
- ocl_calc.mapAndCopy64Bits( dpOclSrcData,mnpOclStartPos,mnpOclEndPos,nSrcDataSize,mnRowSize );
- ocl_calc.oclHostFormulaStatistics64Bits( mcHostName, dpResult, mnRowSize );
+ if (!ocl_calc.mapAndCopy64Bits(dpOclSrcData, mnpOclStartPos, mnpOclEndPos, nSrcDataSize, mnRowSize))
+ return false;
+ if (!ocl_calc.oclHostFormulaStatistics64Bits(mcHostName, dpResult, mnRowSize))
+ return false;
}
else
{
- ocl_calc.createFormulaBuf32Bits( nSrcDataSize, mnPositonLen );
- ocl_calc.mapAndCopy32Bits( dpOclSrcData, mnpOclStartPos, mnpOclEndPos, nSrcDataSize, mnRowSize);
- ocl_calc.oclHostFormulaStatistics32Bits( mcHostName, dpResult, mnRowSize );
+ if (!ocl_calc.createFormulaBuf32Bits(nSrcDataSize, mnPositonLen))
+ return false;
+ if (!ocl_calc.mapAndCopy32Bits(dpOclSrcData, mnpOclStartPos, mnpOclEndPos, nSrcDataSize, mnRowSize))
+ return false;
+ if (!ocl_calc.oclHostFormulaStatistics32Bits(mcHostName, dpResult, mnRowSize))
+ return false;
}
}
}
@@ -359,16 +371,21 @@ bool FormulaGroupInterpreterOpenCL::chooseFunction( OclCalc &ocl_calc, double *&
{
if ( ocl_calc.gpuEnv.mnKhrFp64Flag == 1 || ocl_calc.gpuEnv.mnAmdFp64Flag == 1 )
{
- ocl_calc.createArithmeticOptBuf64Bits( mnRowSize );
- ocl_calc.mapAndCopy64Bits(dpLeftData,dpRightData,mnRowSize);
+ if (!ocl_calc.createArithmeticOptBuf64Bits(mnRowSize))
+ return false;
+ if (!ocl_calc.mapAndCopy64Bits(dpLeftData, dpRightData, mnRowSize))
+ return false;
if (!ocl_calc.oclHostArithmeticOperator64Bits(mcHostName, dpResult, mnRowSize))
return false;
}
else
{
- ocl_calc.createArithmeticOptBuf32Bits( mnRowSize );
- ocl_calc.mapAndCopy32Bits(dpLeftData,dpRightData,mnRowSize);
- ocl_calc.oclHostArithmeticOperator32Bits( mcHostName,dpResult,mnRowSize );
+ if (!ocl_calc.createArithmeticOptBuf32Bits(mnRowSize))
+ return false;
+ if (!ocl_calc.mapAndCopy32Bits(dpLeftData, dpRightData, mnRowSize))
+ return false;
+ if (!ocl_calc.oclHostArithmeticOperator32Bits(mcHostName, dpResult, mnRowSize))
+ return false;
}
}
}
@@ -396,15 +413,21 @@ bool FormulaGroupInterpreterOpenCL::chooseFunction( OclCalc &ocl_calc, double *&
{
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 );
+ if (!ocl_calc.createMoreColArithmeticBuf64Bits(j * mnRowSize, mnOperatorCount))
+ return false;
+ if (!ocl_calc.mapAndCopyMoreColArithmetic64Bits(dpMoreColArithmetic, mnRowSize * j, mnOperatorGroup, mnOperatorCount))
+ return false;
+ if (!ocl_calc.oclMoreColHostArithmeticOperator64Bits(mnRowSize, mnOperatorCount, dpResult, mnRowSize))
+ return false;
}
else
{
- ocl_calc.createMoreColArithmeticBuf32Bits( j* mnRowSize, mnOperatorCount );
- ocl_calc.mapAndCopyMoreColArithmetic32Bits(dpMoreColArithmetic, mnRowSize * j, mnOperatorGroup, mnOperatorCount);
- ocl_calc.oclMoreColHostArithmeticOperator32Bits( mnRowSize, mnOperatorCount, dpResult, mnRowSize );
+ if (!ocl_calc.createMoreColArithmeticBuf32Bits(j* mnRowSize, mnOperatorCount))
+ return false;
+ if (!ocl_calc.mapAndCopyMoreColArithmetic32Bits(dpMoreColArithmetic, mnRowSize * j, mnOperatorGroup, mnOperatorCount))
+ return false;
+ if (!ocl_calc.oclMoreColHostArithmeticOperator32Bits(mnRowSize, mnOperatorCount, dpResult, mnRowSize))
+ return false;
}
}
}
commit d6c03527172bc0ff13c46de6ecc99628508bbf78
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 16:10:35 2013 -0400
Wrong config option name.
Change-Id: Ibbd2d9ed7ee20b8cc911bb5f3247bade617b7c74
diff --git a/sc/source/core/tool/formulaopt.cxx b/sc/source/core/tool/formulaopt.cxx
index 4c926e9..f649927 100644
--- a/sc/source/core/tool/formulaopt.cxx
+++ b/sc/source/core/tool/formulaopt.cxx
@@ -216,7 +216,7 @@ Sequence<OUString> ScFormulaCfg::GetPropertyNames()
"Load/ODFRecalcMode", // SCFORMULAOPT_ODF_RECALC
"Calculation/OpenCL", // SCFORMULAOPT_OPENCL_ENABLED
"Calculation/OpenCLAutoSelect", // SCFORMULAOPT_OPENCL_AUTOSELECT
- "Calculation/OpenCLDevice" // SCFORMULAOPT_OPENCL_DEVICE
+ "Calculation/OpenCLAutoDevice" // SCFORMULAOPT_OPENCL_DEVICE
};
Sequence<OUString> aNames(SCFORMULAOPT_COUNT);
OUString* pNames = aNames.getArray();
commit f46c1196bef3fced26712e8f67e2debd01b4b55e
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 15:56:38 2013 -0400
Pre-compile OpenCL kernels during xlsx file load as well.
Change-Id: Id31190b2a9eb496f146fbbeaaf8a223d46492a99
diff --git a/sc/source/filter/inc/workbookhelper.hxx b/sc/source/filter/inc/workbookhelper.hxx
index e540ba0..9ea2353 100644
--- a/sc/source/filter/inc/workbookhelper.hxx
+++ b/sc/source/filter/inc/workbookhelper.hxx
@@ -275,6 +275,8 @@ public:
/** Returns the codec helper that stores the encoder/decoder object. */
BiffCodecHelper& getCodecHelper() const;
+ void compileOpenCLKernels();
+
private:
WorkbookGlobals& mrBookGlob;
};
diff --git a/sc/source/filter/oox/formulabuffer.cxx b/sc/source/filter/oox/formulabuffer.cxx
index 21e383f..cc8a50a 100644
--- a/sc/source/filter/oox/formulabuffer.cxx
+++ b/sc/source/filter/oox/formulabuffer.cxx
@@ -65,6 +65,7 @@ void FormulaBuffer::finalizeImport()
FormulaDataMap::iterator cellIt = maCellFormulas.find( nTab );
if ( cellIt != maCellFormulas.end() )
{
+ compileOpenCLKernels();
applyCellFormulas( cellIt->second );
}
diff --git a/sc/source/filter/oox/workbookhelper.cxx b/sc/source/filter/oox/workbookhelper.cxx
index 7a9a58f..9ae6fea 100644
--- a/sc/source/filter/oox/workbookhelper.cxx
+++ b/sc/source/filter/oox/workbookhelper.cxx
@@ -71,6 +71,8 @@
#include "dbdata.hxx"
#include "datauno.hxx"
#include "globalnames.hxx"
+#include "clkernelthread.hxx"
+#include "rtl/ref.hxx"
#include "formulabuffer.hxx"
#include "vcl/mapmod.hxx"
@@ -223,6 +225,8 @@ public:
/** Returns the codec helper that stores the encoder/decoder object. */
inline BiffCodecHelper& getCodecHelper() { return *mxCodecHelper; }
+ void compileOpenCLKernels();
+
private:
/** Initializes some basic members and sets needed document properties. */
void initialize( bool bWorkbookFile );
@@ -253,6 +257,8 @@ private:
typedef ::std::auto_ptr< PageSettingsConverter > PageSettConvPtr;
typedef ::std::auto_ptr< BiffCodecHelper > BiffCodecHelperPtr;
+ rtl::Reference<sc::CLBuildKernelThread> mxCLKernelThread;
+
OUString maCellStyles; /// Style family name for cell styles.
OUString maPageStyles; /// Style family name for page styles.
OUString maCellStyleServ; /// Service name for a cell style.
@@ -504,6 +510,15 @@ void WorkbookGlobals::useInternalChartDataTable( bool bInternal )
// private --------------------------------------------------------------------
+void WorkbookGlobals::compileOpenCLKernels()
+{
+ if (mxCLKernelThread.is())
+ return;
+
+ mxCLKernelThread.set(new sc::CLBuildKernelThread);
+ mxCLKernelThread->launch();
+}
+
void WorkbookGlobals::initialize( bool bWorkbookFile )
{
maCellStyles = "CellStyles";
@@ -634,6 +649,9 @@ void WorkbookGlobals::finalize()
//ScDocShell::AfterXMLLoading() for ods
getScDocument().SetInsertingFromOtherDoc(false);
getScDocument().RebuildFormulaGroups();
+
+ if (mxCLKernelThread.is())
+ mxCLKernelThread->join();
}
}
@@ -949,6 +967,11 @@ BiffCodecHelper& WorkbookHelper::getCodecHelper() const
return mrBookGlob.getCodecHelper();
}
+void WorkbookHelper::compileOpenCLKernels()
+{
+ mrBookGlob.compileOpenCLKernels();
+}
+
// ============================================================================
} // namespace xls
commit 2facb76c2a2c0edd6e694378c96302b39d754dcd
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 15:18:01 2013 -0400
Compiler warning.
Change-Id: Ibf3d21d2666b5c2504804141200ea0fcf9897565
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 7a8b205..2b2f8f7 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -56,11 +56,14 @@ namespace {
OString generateHashForSource()
{
- size_t nLength = strlen(kernel_src);
sal_uInt8 pBuffer[RTL_DIGEST_LENGTH_MD5];
+
+#ifndef NDEBUG
+ size_t nLength = strlen(kernel_src);
rtlDigestError aError = rtl_digest_MD5(kernel_src, nLength,
pBuffer, RTL_DIGEST_LENGTH_MD5);
assert(aError == rtl_Digest_E_None);
+#endif
OStringBuffer aBuffer;
const char* pString = "0123456789ABCDEF";
commit 155f878af7e1d3e5cab0acc7a3619c555d4d4d5e
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 15:14:00 2013 -0400
Let's treat oclMatrixSolve equally. No special treatment for this guy.
Change-Id: I79d36ad7c95bf4cc8cd6bb4fd55dcedd5cd70684
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index ce0f662..7a8b205 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -125,7 +125,9 @@ const char* pKernelNames[] = {
"oclMaxDiv",
"oclAverageDiv"
"oclMinDiv",
- "oclSub"
+ "oclSub",
+
+ "oclMatrixSolve"
};
}
@@ -2341,20 +2343,22 @@ bool OclCalc::oclHostMatrixInverse64Bits( const char* aKernelName, double *dpOcl
}
clStatus = clEnqueueUnmapMemObject( kEnv.mpkCmdQueue, mpClmemLeftData, dpOclMatrixSrc, 0, NULL, NULL );
CHECK_OPENCL( clStatus, "clEnqueueUnmapMemObject" );
- cl_kernel kernel_solve = clCreateKernel( kEnv.mpkProgram, "oclMatrixSolve", &clStatus );
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- clStatus = clSetKernelArg( kernel_solve, 0, sizeof(cl_mem), (void *)&mpClmemLeftData );
+ Kernel* pKernelMatrix = fetchKernel("oclMatrixSolve");
+ if (!pKernelMatrix)
+ return false;
+
+ clStatus = clSetKernelArg(pKernel->mpKernel, 0, sizeof(cl_mem), (void *)&mpClmemLeftData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 1, sizeof(cl_mem), (void *)&mpClmemRightData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 1, sizeof(cl_mem), (void *)&mpClmemRightData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 2, sizeof(cl_mem), (void *)&clpPData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 2, sizeof(cl_mem), (void *)&clpPData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 3, sizeof(cl_mem), (void *)&clpYData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 3, sizeof(cl_mem), (void *)&clpYData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 4, sizeof(cl_mem), (void *)&clpNData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 4, sizeof(cl_mem), (void *)&clpNData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clEnqueueNDRangeKernel( kEnv.mpkCmdQueue, kernel_solve, 1, NULL, global_work_size, NULL, 0, NULL, NULL );
+ clStatus = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, pKernel->mpKernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
CHECK_OPENCL( clStatus, "clEnqueueNDRangeKernel" );
clFinish( kEnv.mpkCmdQueue );
for ( uint i = 0; i < nDim; i++ )
@@ -2370,8 +2374,6 @@ bool OclCalc::oclHostMatrixInverse64Bits( const char* aKernelName, double *dpOcl
clStatus = clReleaseMemObject( mpClmemRightData );
CHECK_OPENCL( clStatus, "clReleaseMemObject" );
mpClmemRightData = NULL;
- clStatus = clReleaseKernel( kernel_solve );
- CHECK_OPENCL( clStatus, "clReleaseKernel" );
clStatus = clReleaseMemObject( clpPData );
CHECK_OPENCL( clStatus, "clReleaseKernel" );
clStatus = clReleaseMemObject( clpYData );
@@ -2453,20 +2455,23 @@ bool OclCalc::oclHostMatrixInverse32Bits( const char* aKernelName, float *fpOclM
clStatus = clEnqueueUnmapMemObject( kEnv.mpkCmdQueue, mpClmemLeftData, fpOclMatrixSrc, 0, NULL, NULL );
CHECK_OPENCL( clStatus, "clEnqueueUnmapMemObject" );
- cl_kernel kernel_solve = clCreateKernel( kEnv.mpkProgram, "oclMatrixSolve", &clStatus );
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- clStatus = clSetKernelArg( kernel_solve, 0, sizeof(cl_mem), (void *)&mpClmemLeftData );
+ Kernel* pKernelMatrix = fetchKernel("oclMatrixSolve");
+ if (!pKernelMatrix)
+ return false;
+
+ clStatus = clSetKernelArg(pKernel->mpKernel, 0, sizeof(cl_mem), (void *)&mpClmemLeftData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 1, sizeof(cl_mem), (void *)&mpClmemRightData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 1, sizeof(cl_mem), (void *)&mpClmemRightData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 2, sizeof(cl_mem), (void *)&clpPData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 2, sizeof(cl_mem), (void *)&clpPData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 3, sizeof(cl_mem), (void *)&clpYData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 3, sizeof(cl_mem), (void *)&clpYData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clSetKernelArg( kernel_solve, 4, sizeof(cl_mem), (void *)&clpNData );
+ clStatus = clSetKernelArg(pKernel->mpKernel, 4, sizeof(cl_mem), (void *)&clpNData);
CHECK_OPENCL( clStatus, "clSetKernelArg" );
- clStatus = clEnqueueNDRangeKernel( kEnv.mpkCmdQueue, kernel_solve, 1, NULL, global_work_size, NULL, 0, NULL, NULL );
+ clStatus = clEnqueueNDRangeKernel(
+ kEnv.mpkCmdQueue, pKernel->mpKernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
CHECK_OPENCL( clStatus, "clEnqueueNDRangeKernel" );
clFinish( kEnv.mpkCmdQueue );
for ( uint i = 0; i < nDim; i++ )
@@ -2482,8 +2487,6 @@ bool OclCalc::oclHostMatrixInverse32Bits( const char* aKernelName, float *fpOclM
clStatus = clReleaseMemObject( mpClmemRightData );
CHECK_OPENCL( clStatus, "clReleaseMemObject" );
mpClmemRightData = NULL;
- clStatus = clReleaseKernel( kernel_solve );
- CHECK_OPENCL( clStatus, "clReleaseKernel" );
clStatus = clReleaseMemObject( clpPData );
CHECK_OPENCL( clStatus, "clReleaseKernel" );
clStatus = clReleaseMemObject( clpYData );
commit 2a7589f95877f6c2b2faadd79040fd514ceaaf02
Author: Kohei Yoshida <kohei.yoshida at collabora.com>
Date: Tue Sep 17 14:59:49 2013 -0400
Compile kernel when fetching the Kernel instance.
To make the code a bit cleaner.
Change-Id: Id129cea834e950e422e55e6c2504c1f88c5dbeab
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index c468669..ce0f662 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -152,15 +152,25 @@ int OpenclDevice::setKernelEnv( KernelEnv *envInfo )
return 1;
}
-Kernel* OpenclDevice::checkKernelName( const char *kernelName )
+Kernel* OpenclDevice::fetchKernel( const char *kernelName )
{
+ cl_int nStatus;
for (size_t i = 0, n = gpuEnv.maKernels.size(); i < n; ++i)
{
Kernel* pKernel = &gpuEnv.maKernels[i];
if (!strcasecmp(kernelName, pKernel->mpName))
{
printf("found the kernel named %s.\n", kernelName);
- return pKernel;
+ if (!pKernel->mpKernel && gpuEnv.mpArryPrograms[0])
+ {
+ pKernel->mpKernel = clCreateKernel(gpuEnv.mpArryPrograms[0], kernelName, &nStatus);
+ if (nStatus != CL_SUCCESS)
+ pKernel->mpKernel = NULL;
+
+ printf("Kernel named '%s' has been compiled\n", kernelName);
+ }
+
+ return pKernel->mpKernel ? pKernel : NULL;
}
}
@@ -1000,15 +1010,10 @@ bool OclCalc::oclHostArithmeticOperator64Bits( const char* aKernelName, double *
{
cl_int clStatus = 0;
size_t global_work_size[1];
- Kernel* pKernel = checkKernelName(aKernelName);
+ Kernel* pKernel = fetchKernel(aKernelName);
if (!pKernel)
return false;
- if (!pKernel->mpKernel)
- {
- pKernel->mpKernel = clCreateKernel(kEnv.mpkProgram, aKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
clFinish( kEnv.mpkCmdQueue );
cl_mem clResult = clCreateBuffer( kEnv.mpkContext, CL_MEM_READ_WRITE,
nRowSize * sizeof(double), NULL, &clStatus);
@@ -1048,16 +1053,10 @@ bool OclCalc::oclMoreColHostArithmeticOperator64Bits( int nDataSize,int neOpSize
cl_int clStatus = 0;
size_t global_work_size[1];
const char *aKernelName = "oclMoreColArithmeticOperator";
- Kernel* pKernel = checkKernelName(aKernelName);
+ Kernel* pKernel = fetchKernel(aKernelName);
if (!pKernel)
return false;
- if (!pKernel->mpKernel)
- {
- pKernel->mpKernel = clCreateKernel(kEnv.mpkProgram, aKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
-
cl_mem clResult = clCreateBuffer( kEnv.mpkContext, CL_MEM_READ_WRITE, nRowSize * sizeof(double), NULL, &clStatus );
CHECK_OPENCL( clStatus, "clCreateBuffer" );
clStatus = clSetKernelArg(pKernel->mpKernel, 0, sizeof(cl_mem), (void *)&mpClmemLeftData);
@@ -1095,15 +1094,10 @@ bool OclCalc::oclHostArithmeticStash64Bits( const char* aKernelName, const doubl
cl_int clStatus = 0;
size_t global_work_size[1];
setKernelEnv( &kEnv );
- Kernel* pKernel = checkKernelName(aKernelName);
+ Kernel* pKernel = fetchKernel(aKernelName);
if (!pKernel)
return false;
- if (!pKernel->mpKernel)
- {
- pKernel->mpKernel = clCreateKernel( kEnv.mpkProgram, aKernelName, &clStatus );
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
clFinish( kEnv.mpkCmdQueue );
cl_mem clLeftData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR),
@@ -1149,16 +1143,10 @@ bool OclCalc::oclHostFormulaStash64Bits( const char* aKernelName, const double*
cl_int clStatus = 0;
size_t global_work_size[1];
setKernelEnv( &kEnv );
- Kernel* pKernel = checkKernelName(aKernelName);
+ Kernel* pKernel = fetchKernel(aKernelName);
if (!pKernel)
return false;
- if (!pKernel->mpKernel)
- {
- pKernel->mpKernel = clCreateKernel(kEnv.mpkProgram, aKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
-
cl_mem clSrcData = clCreateBuffer( kEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR),
nBufferSize * sizeof(double), (void *)dpSrcData, &clStatus );
CHECK_OPENCL( clStatus, "clCreateBuffer" );
@@ -1205,16 +1193,10 @@ bool OclCalc::oclHostFormulaStatistics64Bits( const char* aKernelName, double *&
{
cl_int clStatus = 0;
size_t global_work_size[1];
- Kernel* pKernel = checkKernelName(aKernelName);
+ Kernel* pKernel = fetchKernel(aKernelName);
if (!pKernel)
return false;
- if (!pKernel->mpKernel)
- {
- pKernel->mpKernel = clCreateKernel(kEnv.mpkProgram, aKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
-
cl_mem outputCl = clCreateBuffer( kEnv.mpkContext, CL_MEM_READ_WRITE, size * sizeof(double), NULL, &clStatus );
CHECK_OPENCL( clStatus, "clCreateBuffer" );
clStatus = clSetKernelArg(pKernel->mpKernel, 0, sizeof(cl_mem),(void *)&mpClmemSrcData);
@@ -1252,18 +1234,12 @@ bool OclCalc::oclHostFormulaStatistics64Bits( const char* aKernelName, double *&
bool OclCalc::oclHostFormulaCount64Bits( uint *npStartPos, uint *npEndPos, double *&dpOutput, int nSize )
{
const char *cpKernelName = "oclFormulaCount";
- Kernel* pKernel = checkKernelName(cpKernelName);
+ Kernel* pKernel = fetchKernel(cpKernelName);
if (!pKernel)
return false;
cl_int clStatus;
- if (!pKernel->mpKernel)
- {
- pKernel->mpKernel = clCreateKernel(kEnv.mpkProgram, cpKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
-
size_t global_work_size[1];
clStatus = clEnqueueUnmapMemObject( kEnv.mpkCmdQueue, mpClmemStartPos, npStartPos, 0, NULL, NULL );
CHECK_OPENCL( clStatus, "clEnqueueUnmapMemObject" );
@@ -1310,15 +1286,10 @@ bool OclCalc::oclHostFormulaSumProduct64Bits( double *dpSumProMergeLfData, doubl
memset(dpOutput,0,nSize);
const char *cpFirstKernelName = "oclSignedMul";
const char *cpSecondKernelName = "oclFormulaSumproduct";
- Kernel* pKernel1 = checkKernelName(cpFirstKernelName);
+ Kernel* pKernel1 = fetchKernel(cpFirstKernelName);
if (!pKernel1)
return false;
- if (!pKernel1->mpKernel)
- {
- pKernel1->mpKernel = clCreateKernel(kEnv.mpkProgram, cpFirstKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
clStatus = clEnqueueUnmapMemObject( kEnv.mpkCmdQueue, mpClmemMergeLfData, dpSumProMergeLfData, 0, NULL, NULL );
CHECK_OPENCL( clStatus, "clEnqueueUnmapMemObject" );
clFinish(kEnv.mpkCmdQueue);
@@ -1348,15 +1319,10 @@ bool OclCalc::oclHostFormulaSumProduct64Bits( double *dpSumProMergeLfData, doubl
clStatus = clReleaseMemObject( mpClmemMergeRtData );
CHECK_OPENCL( clStatus, "clReleaseMemObject" );
- Kernel* pKernel2 = checkKernelName(cpSecondKernelName);
+ Kernel* pKernel2 = fetchKernel(cpSecondKernelName);
if (!pKernel2)
return false;
- if (!pKernel2->mpKernel)
- {
- pKernel2->mpKernel = clCreateKernel(kEnv.mpkProgram, cpSecondKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
cl_mem clpOutput = clCreateBuffer( kEnv.mpkContext, CL_MEM_READ_WRITE, nSize* sizeof(double), NULL, &clStatus );
CHECK_OPENCL( clStatus, "clCreateBuffer" );
cl_uint nMatixSize = nFormulaColSize * nFormulaRowSize;
@@ -1576,16 +1542,10 @@ bool OclCalc::oclHostArithmeticOperator32Bits( const char* aKernelName,double *r
cl_int clStatus = 0;
size_t global_work_size[1];
- Kernel* pKernel = checkKernelName(aKernelName);
+ Kernel* pKernel = fetchKernel(aKernelName);
if (!pKernel)
return false;
- if (!pKernel->mpKernel)
- {
- pKernel->mpKernel = clCreateKernel(kEnv.mpkProgram, aKernelName, &clStatus);
- CHECK_OPENCL( clStatus, "clCreateKernel" );
- }
-
cl_mem clResult = clCreateBuffer( kEnv.mpkContext, CL_MEM_READ_WRITE, nRowSize * sizeof(float), NULL, &clStatus );
CHECK_OPENCL( clStatus, "clCreateBuffer" );
clStatus = clSetKernelArg(pKernel->mpKernel, 0, sizeof(cl_mem), (void *)&mpClmemLeftData);
... etc. - the rest is truncated
More information about the Libreoffice-commits
mailing list