sc/source/core/data/table2.cxx | 62
sc/source/core/opencl/formulagroupcl.cxx | 2305 ++++++++++++++++---------------
2 files changed, 1236 insertions(+), 1131 deletions(-)
New commits:
commit a23801b740d8ec952097a5e2a2feb450bc0e585f
Author: Kohei Yoshida <kohei.yosh...@collabora.com>
Date: Mon Sep 29 19:08:02 2014 -0400
Scope reduction by early bail-out.
Change-Id: Iacbf9f46a1b5683b6de5dd93002a69078af46951
diff --git a/sc/source/core/data/table2.cxx b/sc/source/core/data/table2.cxx
index 401a148..882c481 100644
--- a/sc/source/core/data/table2.cxx
+++ b/sc/source/core/data/table2.cxx
@@ -484,47 +484,47 @@ void ScTable::CopyToClip(
sc::CopyToClipContext& rCxt, SCCOL nCol1, SCROW nRow1, SCCOL nCol2, SCROW
nRow2,
ScTable* pTable )
{
- if (ValidColRow(nCol1, nRow1) && ValidColRow(nCol2, nRow2))
- {
- // copy content
- //local range names need to be copied first for formula cells
- if (!pTable->mpRangeName && mpRangeName)
- pTable->mpRangeName = new ScRangeName(*mpRangeName);
+ if (!ValidColRow(nCol1, nRow1) || !ValidColRow(nCol2, nRow2))
+ return;
- SCCOL i;
+ // copy content
+ //local range names need to be copied first for formula cells
+ if (!pTable->mpRangeName && mpRangeName)
+ pTable->mpRangeName = new ScRangeName(*mpRangeName);
- for ( i = nCol1; i <= nCol2; i++)
- aCol[i].CopyToClip(rCxt, nRow1, nRow2, pTable->aCol[i]); // notes
are handled at column level
+ SCCOL i;
- // copy widths/heights, and only "hidden", "filtered" and "manual"
flags
- // also for all preceding columns/rows, to have valid positions for
drawing objects
+ for ( i = nCol1; i <= nCol2; i++)
+ aCol[i].CopyToClip(rCxt, nRow1, nRow2, pTable->aCol[i]); // notes are
handled at column level
- if (pColWidth && pTable->pColWidth)
- for (i=0; i<=nCol2; i++)
- pTable->pColWidth[i] = pColWidth[i];
+ // copy widths/heights, and only "hidden", "filtered" and "manual" flags
+ // also for all preceding columns/rows, to have valid positions for
drawing objects
- pTable->CopyColHidden(*this, 0, nCol2);
- pTable->CopyColFiltered(*this, 0, nCol2);
- if (pDBDataNoName)
- pTable->SetAnonymousDBData(new ScDBData(*pDBDataNoName));
+ if (pColWidth && pTable->pColWidth)
+ for (i=0; i<=nCol2; i++)
+ pTable->pColWidth[i] = pColWidth[i];
- if (pRowFlags && pTable->pRowFlags && mpRowHeights &&
pTable->mpRowHeights)
- {
- pTable->pRowFlags->CopyFromAnded( *pRowFlags, 0, nRow2,
CR_MANUALSIZE);
- pTable->CopyRowHeight(*this, 0, nRow2, 0);
- }
+ pTable->CopyColHidden(*this, 0, nCol2);
+ pTable->CopyColFiltered(*this, 0, nCol2);
+ if (pDBDataNoName)
+ pTable->SetAnonymousDBData(new ScDBData(*pDBDataNoName));
+
+ if (pRowFlags && pTable->pRowFlags && mpRowHeights && pTable->mpRowHeights)
+ {
+ pTable->pRowFlags->CopyFromAnded( *pRowFlags, 0, nRow2, CR_MANUALSIZE);
+ pTable->CopyRowHeight(*this, 0, nRow2, 0);
+ }
- pTable->CopyRowHidden(*this, 0, nRow2);
- pTable->CopyRowFiltered(*this, 0, nRow2);
+ pTable->CopyRowHidden(*this, 0, nRow2);
+ pTable->CopyRowFiltered(*this, 0, nRow2);
- // ggf. Formeln durch Werte ersetzen
+ // ggf. Formeln durch Werte ersetzen
- if ( IsProtected() )
- for (i = nCol1; i <= nCol2; i++)
- pTable->aCol[i].RemoveProtected(nRow1, nRow2);
+ if ( IsProtected() )
+ for (i = nCol1; i <= nCol2; i++)
+ pTable->aCol[i].RemoveProtected(nRow1, nRow2);
- pTable->mpCondFormatList.reset(new
ScConditionalFormatList(pTable->pDocument, *mpCondFormatList));
- }
+ pTable->mpCondFormatList.reset(new
ScConditionalFormatList(pTable->pDocument, *mpCondFormatList));
}
void ScTable::CopyToClip(
commit cdd2fe037aa23caeecd28d732df03694e9b62e46
Author: Kohei Yoshida <kohei.yosh...@collabora.com>
Date: Thu Oct 2 15:24:46 2014 -0400
Make the coding style consistent with the rest of Calc code.
I just ran SlickEdit's 'beautify' command with some follow-up manual
editing.
Change-Id: I402e782e7d147c4c95ebaa8ffcc40b50a0c565a7
diff --git a/sc/source/core/opencl/formulagroupcl.cxx
b/sc/source/core/opencl/formulagroupcl.cxx
index 0f61972c..d794e42 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -37,7 +37,9 @@
#ifdef WIN32
#ifndef NAN
namespace {
-static const unsigned long __nan[2] = {0xffffffff, 0x7fffffff};
+
+const unsigned long __nan[2] = {0xffffffff, 0x7fffffff};
+
}
#define NAN (*(const double*) __nan)
#endif
@@ -63,27 +65,32 @@ using namespace formula;
namespace sc { namespace opencl {
/// Map the buffer used by an argument and do necessary argument setting
-size_t VectorRef::Marshal(cl_kernel k, int argno, int, cl_program)
+size_t VectorRef::Marshal( cl_kernel k, int argno, int, cl_program )
{
- FormulaToken *ref = mFormulaTree->GetFormulaToken();
- double *pHostBuffer = NULL;
+ FormulaToken* ref = mFormulaTree->GetFormulaToken();
+ double* pHostBuffer = NULL;
size_t szHostBuffer = 0;
- if (ref->GetType() == formula::svSingleVectorRef) {
+ if (ref->GetType() == formula::svSingleVectorRef)
+ {
const formula::SingleVectorRefToken* pSVR =
- static_cast< const formula::SingleVectorRefToken* >(ref);
+ static_cast<const formula::SingleVectorRefToken*>(ref);
pHostBuffer = const_cast<double*>(pSVR->GetArray().mpNumericArray);
szHostBuffer = pSVR->GetArrayLength() * sizeof(double);
#if 0
std::cerr << "Marshal a Single vector of size " <<
pSVR->GetArrayLength();
std::cerr << " at argument "<< argno << "\n";
#endif
- } else if (ref->GetType() == formula::svDoubleVectorRef) {
+ }
+ else if (ref->GetType() == formula::svDoubleVectorRef)
+ {
const formula::DoubleVectorRefToken* pDVR =
- static_cast< const formula::DoubleVectorRefToken* >(ref);
+ static_cast<const formula::DoubleVectorRefToken*>(ref);
pHostBuffer = const_cast<double*>(
- pDVR->GetArrays()[mnIndex].mpNumericArray);
+ pDVR->GetArrays()[mnIndex].mpNumericArray);
szHostBuffer = pDVR->GetArrayLength() * sizeof(double);
- } else {
+ }
+ else
+ {
throw Unhandled();
}
// Obtain cl context
@@ -93,9 +100,9 @@ size_t VectorRef::Marshal(cl_kernel k, int argno, int,
cl_program)
if (pHostBuffer)
{
mpClmem = clCreateBuffer(kEnv.mpkContext,
- (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR,
- szHostBuffer,
- pHostBuffer, &err);
+ (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+ szHostBuffer,
+ pHostBuffer, &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
}
@@ -103,21 +110,21 @@ size_t VectorRef::Marshal(cl_kernel k, int argno, int,
cl_program)
{
if (szHostBuffer == 0)
szHostBuffer = sizeof(double); // a dummy small value
- // Marshal as a buffer of NANs
+ // Marshal as a buffer of NANs
mpClmem = clCreateBuffer(kEnv.mpkContext,
- (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR,
- szHostBuffer, NULL, &err);
+ (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,
+ szHostBuffer, NULL, &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
- double *pNanBuffer = (double*)clEnqueueMapBuffer(
- kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
- szHostBuffer, 0, NULL, NULL, &err);
+ double* pNanBuffer = (double*)clEnqueueMapBuffer(
+ kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
+ szHostBuffer, 0, NULL, NULL, &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
- for (size_t i = 0; i < szHostBuffer/sizeof(double); i++)
+ for (size_t i = 0; i < szHostBuffer / sizeof(double); i++)
pNanBuffer[i] = NAN;
err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem,
- pNanBuffer, 0, NULL, NULL);
+ pNanBuffer, 0, NULL, NULL);
}
err = clSetKernelArg(k, argno, sizeof(cl_mem), (void*)&mpClmem);
@@ -130,48 +137,50 @@ size_t VectorRef::Marshal(cl_kernel k, int argno, int,
cl_program)
/// Currently, only the hash is passed.
/// TBD(IJSUNG): pass also length and the actual string if there is a
/// hash function collision
-class ConstStringArgument: public DynamicKernelArgument
+class ConstStringArgument : public DynamicKernelArgument
{
public:
- ConstStringArgument(const std::string &s,
- FormulaTreeNodeRef ft):
- DynamicKernelArgument(s, ft) {}
+ ConstStringArgument( const std::string& s,
+ FormulaTreeNodeRef ft ) :
+ DynamicKernelArgument(s, ft) { }
/// Generate declaration
- virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << "unsigned " << mSymName;
}
- virtual void GenDeclRef(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << GenSlidingWindowDeclRef(false);
}
- virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenSlidingWindowDecl( std::stringstream& ss ) const
SAL_OVERRIDE
{
GenDecl(ss);
}
- virtual std::string GenSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE
+ virtual std::string GenSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
std::stringstream ss;
if (GetFormulaToken()->GetType() != formula::svString)
throw Unhandled();
- FormulaToken *Tok = GetFormulaToken();
+ FormulaToken* Tok = GetFormulaToken();
ss << Tok->GetString().getString().toAsciiUpperCase().hashCode() <<
"U";
return ss.str();
}
- virtual size_t GetWindowSize(void) const SAL_OVERRIDE
+ virtual size_t GetWindowSize() const SAL_OVERRIDE
{
return 1;
}
/// Pass the 32-bit hash of the string to the kernel
- virtual size_t Marshal(cl_kernel k, int argno, int, cl_program)
SAL_OVERRIDE
+ virtual size_t Marshal( cl_kernel k, int argno, int, cl_program )
SAL_OVERRIDE
{
- FormulaToken *ref = mFormulaTree->GetFormulaToken();
+ FormulaToken* ref = mFormulaTree->GetFormulaToken();
cl_uint hashCode = 0;
if (ref->GetType() == formula::svString)
{
const rtl::OUString s =
ref->GetString().getString().toAsciiUpperCase();
hashCode = s.hashCode();
- } else {
+ }
+ else
+ {
throw Unhandled();
}
// marshaling
@@ -187,44 +196,44 @@ public:
};
/// Arguments that are actually compile-time constants
-class DynamicKernelConstantArgument: public DynamicKernelArgument
+class DynamicKernelConstantArgument : public DynamicKernelArgument
{
public:
- DynamicKernelConstantArgument(const std::string &s,
- FormulaTreeNodeRef ft):
- DynamicKernelArgument(s, ft) {}
+ DynamicKernelConstantArgument( const std::string& s,
+ FormulaTreeNodeRef ft ) :
+ DynamicKernelArgument(s, ft) { }
/// Generate declaration
- virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << "double " << mSymName;
}
- virtual void GenDeclRef(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << mSymName;
}
- virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenSlidingWindowDecl( std::stringstream& ss ) const
SAL_OVERRIDE
{
GenDecl(ss);
}
- virtual std::string GenSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE
+ virtual std::string GenSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
if (GetFormulaToken()->GetType() != formula::svDouble)
throw Unhandled();
return mSymName;
}
- virtual size_t GetWindowSize(void) const SAL_OVERRIDE
+ virtual size_t GetWindowSize() const SAL_OVERRIDE
{
return 1;
}
- double GetDouble(void) const
+ double GetDouble() const
{
- FormulaToken *Tok = GetFormulaToken();
+ FormulaToken* Tok = GetFormulaToken();
if (Tok->GetType() != formula::svDouble)
throw Unhandled();
return Tok->GetDouble();
}
/// Create buffer and pass the buffer to a given kernel
- virtual size_t Marshal(cl_kernel k, int argno, int, cl_program)
SAL_OVERRIDE
+ virtual size_t Marshal( cl_kernel k, int argno, int, cl_program )
SAL_OVERRIDE
{
double tmp = GetDouble();
// Pass the scalar result back to the rest of the formula kernel
@@ -233,38 +242,38 @@ public:
throw OpenCLError(err, __FILE__, __LINE__);
return 1;
}
- virtual cl_mem GetCLBuffer(void) const { return NULL; }
+ virtual cl_mem GetCLBuffer() const { return NULL; }
};
-class DynamicKernelPiArgument: public DynamicKernelArgument
+class DynamicKernelPiArgument : public DynamicKernelArgument
{
public:
- DynamicKernelPiArgument(const std::string &s,
- FormulaTreeNodeRef ft):
- DynamicKernelArgument(s, ft) {}
+ DynamicKernelPiArgument( const std::string& s,
+ FormulaTreeNodeRef ft ) :
+ DynamicKernelArgument(s, ft) { }
/// Generate declaration
- virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << "double " << mSymName;
}
- virtual void GenDeclRef(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << "3.14159265358979";
}
- virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenSlidingWindowDecl( std::stringstream& ss ) const
SAL_OVERRIDE
{
GenDecl(ss);
}
- virtual std::string GenSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE
+ virtual std::string GenSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
return mSymName;
}
- virtual size_t GetWindowSize(void) const SAL_OVERRIDE
+ virtual size_t GetWindowSize() const SAL_OVERRIDE
{
return 1;
}
/// Create buffer and pass the buffer to a given kernel
- virtual size_t Marshal(cl_kernel k, int argno, int, cl_program)
SAL_OVERRIDE
+ virtual size_t Marshal( cl_kernel k, int argno, int, cl_program )
SAL_OVERRIDE
{
double tmp = 0.0;
// Pass the scalar result back to the rest of the formula kernel
@@ -275,30 +284,30 @@ public:
}
};
-class DynamicKernelRandomArgument: public DynamicKernelArgument
+class DynamicKernelRandomArgument : public DynamicKernelArgument
{
public:
- DynamicKernelRandomArgument(const std::string &s,
- FormulaTreeNodeRef ft):
- DynamicKernelArgument(s, ft) {}
+ DynamicKernelRandomArgument( const std::string& s,
+ FormulaTreeNodeRef ft ) :
+ DynamicKernelArgument(s, ft) { }
/// Generate declaration
- virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << "double " << mSymName;
}
- virtual void GenDeclRef(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
ss << mSymName;
}
- virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenSlidingWindowDecl( std::stringstream& ss ) const
SAL_OVERRIDE
{
GenDecl(ss);
}
- virtual std::string GenSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE
+ virtual std::string GenSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
return mSymName + "_Random()";
}
- void GenSlidingWindowFunction(std::stringstream &ss) SAL_OVERRIDE
+ void GenSlidingWindowFunction( std::stringstream& ss ) SAL_OVERRIDE
{
ss << "\ndouble " << mSymName;
ss << "_Random ()\n{\n";
@@ -318,12 +327,12 @@ public:
ss << " return tmp;\n";
ss << "}";
}
- virtual size_t GetWindowSize(void) const SAL_OVERRIDE
+ virtual size_t GetWindowSize() const SAL_OVERRIDE
{
return 1;
}
/// Create buffer and pass the buffer to a given kernel
- virtual size_t Marshal(cl_kernel k, int argno, int, cl_program)
SAL_OVERRIDE
+ virtual size_t Marshal( cl_kernel k, int argno, int, cl_program )
SAL_OVERRIDE
{
double tmp = 0.0;
// Pass the scalar result back to the rest of the formula kernel
@@ -335,62 +344,65 @@ public:
};
/// A vector of strings
-class DynamicKernelStringArgument: public VectorRef
+class DynamicKernelStringArgument : public VectorRef
{
public:
- DynamicKernelStringArgument(const std::string &s,
- FormulaTreeNodeRef ft, int index = 0):
- VectorRef(s, ft, index) {}
+ DynamicKernelStringArgument( const std::string& s,
+ FormulaTreeNodeRef ft, int index = 0 ) :
+ VectorRef(s, ft, index) { }
- virtual void GenSlidingWindowFunction(std::stringstream &) SAL_OVERRIDE {}
+ virtual void GenSlidingWindowFunction( std::stringstream& ) SAL_OVERRIDE {
}
/// Generate declaration
- virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE
{
- ss << "__global unsigned int *"<<mSymName;
+ ss << "__global unsigned int *" << mSymName;
}
- virtual void GenSlidingWindowDecl(std::stringstream& ss) const SAL_OVERRIDE
+ virtual void GenSlidingWindowDecl( std::stringstream& ss ) const
SAL_OVERRIDE
{
DynamicKernelStringArgument::GenDecl(ss);
}
- virtual size_t Marshal(cl_kernel, int, int, cl_program) SAL_OVERRIDE;
+ virtual size_t Marshal( cl_kernel, int, int, cl_program ) SAL_OVERRIDE;
};
/// Marshal a string vector reference
-size_t DynamicKernelStringArgument::Marshal(cl_kernel k, int argno, int,
cl_program)
+size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int,
cl_program )
{
- FormulaToken *ref = mFormulaTree->GetFormulaToken();
+ FormulaToken* ref = mFormulaTree->GetFormulaToken();
// Obtain cl context
KernelEnv kEnv;
OpenclDevice::setKernelEnv(&kEnv);
cl_int err;
formula::VectorRefArray vRef;
size_t nStrings = 0;
- if (ref->GetType() == formula::svSingleVectorRef) {
+ if (ref->GetType() == formula::svSingleVectorRef)
+ {
const formula::SingleVectorRefToken* pSVR =
- static_cast< const formula::SingleVectorRefToken* >(ref);
+ static_cast<const formula::SingleVectorRefToken*>(ref);
nStrings = pSVR->GetArrayLength();
vRef = pSVR->GetArray();
- } else if (ref->GetType() == formula::svDoubleVectorRef) {
+ }
+ else if (ref->GetType() == formula::svDoubleVectorRef)
+ {
const formula::DoubleVectorRefToken* pDVR =
- static_cast< const formula::DoubleVectorRefToken* >(ref);
+ static_cast<const formula::DoubleVectorRefToken*>(ref);
nStrings = pDVR->GetArrayLength();
vRef = pDVR->GetArrays()[mnIndex];
}
size_t szHostBuffer = nStrings * sizeof(cl_int);
- cl_uint *pHashBuffer = NULL;
+ cl_uint* pHashBuffer = NULL;
- if ( vRef.mpStringArray != NULL)
+ if (vRef.mpStringArray != NULL)
{
// Marshal strings. Right now we pass hashes of these string
mpClmem = clCreateBuffer(kEnv.mpkContext,
- (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR,
- szHostBuffer, NULL, &err);
+ (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,
+ szHostBuffer, NULL, &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
pHashBuffer = (cl_uint*)clEnqueueMapBuffer(
- kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
- szHostBuffer, 0, NULL, NULL, &err);
+ kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
+ szHostBuffer, 0, NULL, NULL, &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
@@ -405,30 +417,30 @@ size_t DynamicKernelStringArgument::Marshal(cl_kernel k,
int argno, int, cl_prog
{
pHashBuffer[i] = 0;
}
- }
+ }
}
else
{
if (nStrings == 0)
szHostBuffer = sizeof(cl_int); // a dummy small value
- // Marshal as a buffer of NANs
+ // Marshal as a buffer of NANs
mpClmem = clCreateBuffer(kEnv.mpkContext,
- (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR,
- szHostBuffer, NULL, &err);
+ (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,
+ szHostBuffer, NULL, &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
pHashBuffer = (cl_uint*)clEnqueueMapBuffer(
- kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
- szHostBuffer, 0, NULL, NULL, &err);
+ kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
+ szHostBuffer, 0, NULL, NULL, &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
- for (size_t i = 0; i < szHostBuffer/sizeof(cl_int); i++)
+ for (size_t i = 0; i < szHostBuffer / sizeof(cl_int); i++)
pHashBuffer[i] = 0;
}
err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem,
- pHashBuffer, 0, NULL, NULL);
+ pHashBuffer, 0, NULL, NULL);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
@@ -439,42 +451,42 @@ size_t DynamicKernelStringArgument::Marshal(cl_kernel k,
int argno, int, cl_prog
}
/// A mixed string/numberic vector
-class DynamicKernelMixedArgument: public VectorRef
+class DynamicKernelMixedArgument : public VectorRef
{
public:
- DynamicKernelMixedArgument(const std::string &s,
- FormulaTreeNodeRef ft):
- VectorRef(s, ft), mStringArgument(s+"s", ft) {}
- virtual void GenSlidingWindowDecl(std::stringstream& ss) const SAL_OVERRIDE
+ DynamicKernelMixedArgument( const std::string& s,
+ FormulaTreeNodeRef ft ) :
+ VectorRef(s, ft), mStringArgument(s + "s", ft) { }
+ virtual void GenSlidingWindowDecl( std::stringstream& ss ) const
SAL_OVERRIDE
{
VectorRef::GenSlidingWindowDecl(ss);
ss << ", ";
mStringArgument.GenSlidingWindowDecl(ss);
}
- virtual bool IsMixedArgument() const SAL_OVERRIDE {return true;}
- virtual void GenSlidingWindowFunction(std::stringstream &) SAL_OVERRIDE {}
+ virtual bool IsMixedArgument() const SAL_OVERRIDE { return true;}
+ virtual void GenSlidingWindowFunction( std::stringstream& ) SAL_OVERRIDE {
}
/// Generate declaration
- virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE
{
VectorRef::GenDecl(ss);
ss << ", ";
mStringArgument.GenDecl(ss);
}
- virtual void GenDeclRef(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
VectorRef::GenDeclRef(ss);
ss << ",";
mStringArgument.GenDeclRef(ss);
}
- virtual void GenNumDeclRef(std::stringstream& ss) const SAL_OVERRIDE
+ virtual void GenNumDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
VectorRef::GenSlidingWindowDecl(ss);
}
- virtual void GenStringDeclRef(std::stringstream& ss) const SAL_OVERRIDE
+ virtual void GenStringDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
mStringArgument.GenSlidingWindowDecl(ss);
}
- virtual std::string GenSlidingWindowDeclRef(bool nested) const SAL_OVERRIDE
+ virtual std::string GenSlidingWindowDeclRef( bool nested ) const
SAL_OVERRIDE
{
std::stringstream ss;
ss << "(!isNan(" << VectorRef::GenSlidingWindowDeclRef();
@@ -483,24 +495,25 @@ public:
ss << ")";
return ss.str();
}
- virtual std::string GenDoubleSlidingWindowDeclRef(bool=false) const
SAL_OVERRIDE
+ virtual std::string GenDoubleSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
std::stringstream ss;
ss << VectorRef::GenSlidingWindowDeclRef();
return ss.str();
}
- virtual std::string GenStringSlidingWindowDeclRef(bool=false) const
SAL_OVERRIDE
+ virtual std::string GenStringSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
std::stringstream ss;
ss << mStringArgument.GenSlidingWindowDeclRef();
return ss.str();
}
- virtual size_t Marshal(cl_kernel k, int argno, int vw, cl_program p)
SAL_OVERRIDE
+ virtual size_t Marshal( cl_kernel k, int argno, int vw, cl_program p )
SAL_OVERRIDE
{
int i = VectorRef::Marshal(k, argno, vw, p);
- i += mStringArgument.Marshal(k, argno+i, vw, p);
+ i += mStringArgument.Marshal(k, argno + i, vw, p);
return i;
}
+
protected:
DynamicKernelStringArgument mStringArgument;
};
@@ -513,40 +526,41 @@ class OpAverage; // Forward Declaration
class OpMin; // Forward Declaration
class OpMax; // Forward Declaration
class OpCount; // Forward Declaration
+
template<class Base>
-class DynamicKernelSlidingArgument: public Base
+class DynamicKernelSlidingArgument : public Base
{
public:
- DynamicKernelSlidingArgument(const std::string &s,
- FormulaTreeNodeRef ft, boost::shared_ptr<SlidingFunctionBase> &CodeGen,
- int index=0):
+ DynamicKernelSlidingArgument( const std::string& s,
+ FormulaTreeNodeRef ft, boost::shared_ptr<SlidingFunctionBase>& CodeGen,
+ int index = 0 ) :
Base(s, ft, index), mpCodeGen(CodeGen), mpClmem2(NULL)
{
- FormulaToken *t = ft->GetFormulaToken();
+ FormulaToken* t = ft->GetFormulaToken();
if (t->GetType() != formula::svDoubleVectorRef)
throw Unhandled();
- mpDVR = static_cast<const formula::DoubleVectorRefToken *>(t);
+ mpDVR = static_cast<const formula::DoubleVectorRefToken*>(t);
bIsStartFixed = mpDVR->IsStartFixed();
bIsEndFixed = mpDVR->IsEndFixed();
}
// Should only be called by SumIfs. Yikes!
- virtual bool NeedParallelReduction(void) const
+ virtual bool NeedParallelReduction() const
{
assert(dynamic_cast<OpSumIfs*>(mpCodeGen.get()));
- return GetWindowSize()> 100 &&
- ( (GetStartFixed() && GetEndFixed()) ||
- (!GetStartFixed() && !GetEndFixed()) ) ;
+ return GetWindowSize() > 100 &&
+ ((GetStartFixed() && GetEndFixed()) ||
+ (!GetStartFixed() && !GetEndFixed()));
}
- virtual void GenSlidingWindowFunction(std::stringstream &) {}
+ virtual void GenSlidingWindowFunction( std::stringstream& ) { }
- virtual std::string GenSlidingWindowDeclRef(bool nested=false) const
+ virtual std::string GenSlidingWindowDeclRef( bool nested = false ) const
{
size_t nArrayLength = mpDVR->GetArrayLength();
std::stringstream ss;
if (!bIsStartFixed && !bIsEndFixed)
{
if (nested)
- ss << "((i+gid0) <" << nArrayLength <<"?";
+ ss << "((i+gid0) <" << nArrayLength << "?";
ss << Base::GetName() << "[i + gid0]";
if (nested)
ss << ":NAN)";
@@ -554,7 +568,7 @@ public:
else
{
if (nested)
- ss << "(i <" << nArrayLength <<"?";
+ ss << "(i <" << nArrayLength << "?";
ss << Base::GetName() << "[i]";
if (nested)
ss << ":NAN)";
@@ -563,7 +577,7 @@ public:
}
/// Controls how the elements in the DoubleVectorRef are traversed
virtual size_t GenReductionLoopHeader(
- std::stringstream &ss, bool &needBody)
+ std::stringstream& ss, bool& needBody )
{
assert(mpDVR);
size_t nCurWindowSize = mpDVR->GetRefRowSize();
@@ -585,61 +599,69 @@ public:
ss << "gid0; i < " << mpDVR->GetArrayLength();
ss << " && i < " << nCurWindowSize << "; i++){\n\t\t";
#else
- ss << "gid0; i < "<< nCurWindowSize << "; i++)\n\t\t";
+ ss << "gid0; i < " << nCurWindowSize << "; i++)\n\t\t";
#endif
}
else if (bIsStartFixed && !bIsEndFixed)
{
#ifdef ISNAN
ss << "0; i < " << mpDVR->GetArrayLength();
- ss << " && i < gid0+"<< nCurWindowSize << "; i++){\n\t\t";
+ ss << " && i < gid0+" << nCurWindowSize << "; i++){\n\t\t";
#else
- ss << "0; i < gid0+"<< nCurWindowSize << "; i++)\n\t\t";
+ ss << "0; i < gid0+" << nCurWindowSize << "; i++)\n\t\t";
#endif
}
else if (!bIsStartFixed && !bIsEndFixed)
{
#ifdef ISNAN
ss << "0; i + gid0 < " << mpDVR->GetArrayLength();
- ss << " && i < "<< nCurWindowSize << "; i++){\n\t\t";
+ ss << " && i < " << nCurWindowSize << "; i++){\n\t\t";
#else
- ss << "0; i < "<< nCurWindowSize << "; i++)\n\t\t";
+ ss << "0; i < " << nCurWindowSize << "; i++)\n\t\t";
#endif
}
else
{
unsigned limit =
std::min(mpDVR->GetArrayLength(), nCurWindowSize);
- ss << "0; i < "<< limit << "; i++){\n\t\t";
+ ss << "0; i < " << limit << "; i++){\n\t\t";
}
-return nCurWindowSize;
+ return nCurWindowSize;
#endif
#ifdef UNROLLING_FACTOR
{
- if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) {
+ if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
+ {
ss << "for (int i = ";
ss << "gid0; i < " << mpDVR->GetArrayLength();
ss << " && i < " << nCurWindowSize << "; i++){\n\t\t";
needBody = true;
return nCurWindowSize;
- } else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) {
+ }
+ else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ {
ss << "for (int i = ";
ss << "0; i < " << mpDVR->GetArrayLength();
- ss << " && i < gid0+"<< nCurWindowSize << "; i++){\n\t\t";
+ ss << " && i < gid0+" << nCurWindowSize << "; i++){\n\t\t";
needBody = true;
return nCurWindowSize;
- } else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()){
+ }
+ else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ {
ss << "tmpBottom = " << mpCodeGen->GetBottom() << ";\n\t";
ss << "{int i;\n\t";
- std::stringstream temp1,temp2;
+ std::stringstream temp1, temp2;
int outLoopSize = UNROLLING_FACTOR;
- if ( nCurWindowSize/outLoopSize != 0){
- ss << "for(int outLoop=0; outLoop<" <<
nCurWindowSize/outLoopSize<< "; outLoop++){\n\t";
- for(int count=0; count < outLoopSize; count++){
- ss << "i = outLoop*"<<outLoopSize<<"+"<<count<<";\n\t";
- if(count==0){
- temp1 << "if(i + gid0 < "
<<mpDVR->GetArrayLength();
+ if (nCurWindowSize / outLoopSize != 0)
+ {
+ ss << "for(int outLoop=0; outLoop<" << nCurWindowSize /
outLoopSize << "; outLoop++){\n\t";
+ for (int count = 0; count < outLoopSize; count++)
+ {
+ ss << "i = outLoop*" << outLoopSize << "+" << count <<
";\n\t";
+ if (count == 0)
+ {
+ temp1 << "if(i + gid0 < " <<
mpDVR->GetArrayLength();
temp1 << "){\n\t\t";
temp1 << "tmp = legalize(";
temp1 <<
mpCodeGen->Gen2(GenSlidingWindowDeclRef(), "tmp");
@@ -651,9 +673,11 @@ return nCurWindowSize;
ss << "}\n\t";
}
// The residual of mod outLoopSize
- for(unsigned int count=nCurWindowSize/outLoopSize*outLoopSize;
count < nCurWindowSize; count++){
- ss << "i = "<<count<<";\n\t";
- if(count==nCurWindowSize/outLoopSize*outLoopSize){
+ for (unsigned int count = nCurWindowSize / outLoopSize *
outLoopSize; count < nCurWindowSize; count++)
+ {
+ ss << "i = " << count << ";\n\t";
+ if (count == nCurWindowSize / outLoopSize * outLoopSize)
+ {
temp2 << "if(i + gid0 < " << mpDVR->GetArrayLength();
temp2 << "){\n\t\t";
temp2 << "tmp = legalize(";
@@ -668,17 +692,21 @@ return nCurWindowSize;
return nCurWindowSize;
}
// (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
- else {
+ else
+ {
ss << "//else situation \n\t";
ss << "tmpBottom = " << mpCodeGen->GetBottom() << ";\n\t";
ss << "{int i;\n\t";
- std::stringstream temp1,temp2;
+ std::stringstream temp1, temp2;
int outLoopSize = UNROLLING_FACTOR;
- if (nCurWindowSize/outLoopSize != 0){
- ss << "for(int outLoop=0; outLoop<" <<
nCurWindowSize/outLoopSize<< "; outLoop++){\n\t";
- for(int count=0; count < outLoopSize; count++){
- ss << "i = outLoop*"<<outLoopSize<<"+"<<count<<";\n\t";
- if(count==0){
+ if (nCurWindowSize / outLoopSize != 0)
+ {
+ ss << "for(int outLoop=0; outLoop<" << nCurWindowSize /
outLoopSize << "; outLoop++){\n\t";
+ for (int count = 0; count < outLoopSize; count++)
+ {
+ ss << "i = outLoop*" << outLoopSize << "+" << count <<
";\n\t";
+ if (count == 0)
+ {
temp1 << "tmp = legalize(";
temp1 <<
mpCodeGen->Gen2(GenSlidingWindowDeclRef(), "tmp");
temp1 << ", tmp);\n\t\t\t";
@@ -688,9 +716,11 @@ return nCurWindowSize;
ss << "}\n\t";
}
// The residual of mod outLoopSize
- for(unsigned int count=nCurWindowSize/outLoopSize*outLoopSize;
count < nCurWindowSize; count++){
- ss << "i = "<<count<<";\n\t";
- if(count==nCurWindowSize/outLoopSize*outLoopSize){
+ for (unsigned int count = nCurWindowSize / outLoopSize *
outLoopSize; count < nCurWindowSize; count++)
+ {
+ ss << "i = " << count << ";\n\t";
+ if (count == nCurWindowSize / outLoopSize * outLoopSize)
+ {
temp2 << "tmp = legalize(";
temp2 << mpCodeGen->Gen2(GenSlidingWindowDeclRef(),
"tmp");
temp2 << ", tmp);\n\t\t\t";
@@ -703,7 +733,7 @@ return nCurWindowSize;
}
}
#endif
-}
+ }
~DynamicKernelSlidingArgument()
{
if (mpClmem2)
@@ -713,17 +743,17 @@ return nCurWindowSize;
}
}
- size_t GetArrayLength(void) const {return mpDVR->GetArrayLength(); }
+ size_t GetArrayLength() const { return mpDVR->GetArrayLength(); }
- size_t GetWindowSize(void) const {return mpDVR->GetRefRowSize(); }
+ size_t GetWindowSize() const { return mpDVR->GetRefRowSize(); }
- size_t GetStartFixed(void) const {return bIsStartFixed; }
+ size_t GetStartFixed() const { return bIsStartFixed; }
- size_t GetEndFixed(void) const {return bIsEndFixed; }
+ size_t GetEndFixed() const { return bIsEndFixed; }
protected:
bool bIsStartFixed, bIsEndFixed;
- const formula::DoubleVectorRefToken *mpDVR;
+ const formula::DoubleVectorRefToken* mpDVR;
// from parent nodes
boost::shared_ptr<SlidingFunctionBase> mpCodeGen;
// controls whether to invoke the reduction kernel during marshaling or not
@@ -734,33 +764,33 @@ protected:
class DynamicKernelMixedSlidingArgument : public VectorRef
{
public:
- DynamicKernelMixedSlidingArgument(const std::string &s,
- FormulaTreeNodeRef ft, boost::shared_ptr<SlidingFunctionBase> &CodeGen,
- int index = 0):
+ DynamicKernelMixedSlidingArgument( const std::string& s,
+ FormulaTreeNodeRef ft, boost::shared_ptr<SlidingFunctionBase>& CodeGen,
+ int index = 0 ) :
VectorRef(s, ft),
mDoubleArgument(s, ft, CodeGen, index),
- mStringArgument(s+"s", ft, CodeGen, index) {}
- virtual void GenSlidingWindowDecl(std::stringstream& ss) const SAL_OVERRIDE
+ mStringArgument(s + "s", ft, CodeGen, index) { }
+ virtual void GenSlidingWindowDecl( std::stringstream& ss ) const
SAL_OVERRIDE
{
mDoubleArgument.GenSlidingWindowDecl(ss);
ss << ", ";
mStringArgument.GenSlidingWindowDecl(ss);
}
- virtual void GenSlidingWindowFunction(std::stringstream &) SAL_OVERRIDE {}
+ virtual void GenSlidingWindowFunction( std::stringstream& ) SAL_OVERRIDE {
}
/// Generate declaration
- virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE
{
mDoubleArgument.GenDecl(ss);
ss << ", ";
mStringArgument.GenDecl(ss);
}
- virtual void GenDeclRef(std::stringstream &ss) const SAL_OVERRIDE
+ virtual void GenDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
mDoubleArgument.GenDeclRef(ss);
ss << ",";
mStringArgument.GenDeclRef(ss);
}
- virtual std::string GenSlidingWindowDeclRef(bool nested) const SAL_OVERRIDE
+ virtual std::string GenSlidingWindowDeclRef( bool nested ) const
SAL_OVERRIDE
{
std::stringstream ss;
ss << "(!isNan(" << mDoubleArgument.GenSlidingWindowDeclRef();
@@ -769,61 +799,66 @@ public:
ss << ")";
return ss.str();
}
- virtual bool IsMixedArgument() const SAL_OVERRIDE {return true;}
- virtual std::string GenDoubleSlidingWindowDeclRef(bool=false) const
SAL_OVERRIDE
+ virtual bool IsMixedArgument() const SAL_OVERRIDE { return true;}
+ virtual std::string GenDoubleSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
std::stringstream ss;
ss << mDoubleArgument.GenSlidingWindowDeclRef();
return ss.str();
}
- virtual std::string GenStringSlidingWindowDeclRef(bool=false) const
SAL_OVERRIDE
+ virtual std::string GenStringSlidingWindowDeclRef( bool = false ) const
SAL_OVERRIDE
{
std::stringstream ss;
ss << mStringArgument.GenSlidingWindowDeclRef();
return ss.str();
}
- virtual void GenNumDeclRef(std::stringstream& ss) const SAL_OVERRIDE
+ virtual void GenNumDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
- mDoubleArgument.GenDeclRef(ss);
+ mDoubleArgument.GenDeclRef(ss);
}
- virtual void GenStringDeclRef(std::stringstream& ss) const SAL_OVERRIDE
+ virtual void GenStringDeclRef( std::stringstream& ss ) const SAL_OVERRIDE
{
mStringArgument.GenDeclRef(ss);
}
- virtual size_t Marshal(cl_kernel k, int argno, int vw, cl_program p)
SAL_OVERRIDE
+ virtual size_t Marshal( cl_kernel k, int argno, int vw, cl_program p )
SAL_OVERRIDE
{
int i = mDoubleArgument.Marshal(k, argno, vw, p);
i += mStringArgument.Marshal(k, argno + i, vw, p);
return i;
}
+
protected:
DynamicKernelSlidingArgument<VectorRef> mDoubleArgument;
DynamicKernelSlidingArgument<DynamicKernelStringArgument> mStringArgument;
};
+
/// Holds the symbol table for a given dynamic kernel
-class SymbolTable {
+class SymbolTable
+{
public:
- typedef std::map<const formula::FormulaToken *,
- boost::shared_ptr<DynamicKernelArgument> > ArgumentMap;
+ typedef std::map<const formula::FormulaToken*,
+ boost::shared_ptr<DynamicKernelArgument> > ArgumentMap;
// This avoids instability caused by using pointer as the key type
- typedef std::list< boost::shared_ptr<DynamicKernelArgument> > ArgumentList;
- SymbolTable(void):mCurId(0) {}
- template <class T>
- const DynamicKernelArgument *DeclRefArg(FormulaTreeNodeRef,
SlidingFunctionBase* pCodeGen);
+ typedef std::list<boost::shared_ptr<DynamicKernelArgument> > ArgumentList;
+ SymbolTable() : mCurId(0) { }
+ template<class T>
+ const DynamicKernelArgument* DeclRefArg( FormulaTreeNodeRef,
SlidingFunctionBase* pCodeGen );
/// Used to generate sliding window helpers
- void DumpSlidingWindowFunctions(std::stringstream &ss)
+ void DumpSlidingWindowFunctions( std::stringstream& ss )
{
- for(ArgumentList::iterator it = mParams.begin(), e= mParams.end();
it!=e;
- ++it) {
+ for (ArgumentList::iterator it = mParams.begin(), e = mParams.end();
it != e;
+ ++it)
+ {
(*it)->GenSlidingWindowFunction(ss);
ss << "\n";
}
}
/// Memory mapping from host to device and pass buffers to the given
kernel as
/// arguments
- void Marshal(cl_kernel, int, cl_program);
+ void Marshal( cl_kernel, int, cl_program );
// number of result items.
static int nR;
+
private:
unsigned int mCurId;
ArgumentMap mSymbols;
@@ -831,212 +866,215 @@ private:
};
int SymbolTable::nR = 0;
-void SymbolTable::Marshal(cl_kernel k, int nVectorWidth, cl_program pProgram)
+void SymbolTable::Marshal( cl_kernel k, int nVectorWidth, cl_program pProgram )
{
int i = 1; //The first argument is reserved for results
- for(ArgumentList::iterator it = mParams.begin(), e= mParams.end(); it!=e;
- ++it) {
- i+=(*it)->Marshal(k, i, nVectorWidth, pProgram);
+ for (ArgumentList::iterator it = mParams.begin(), e = mParams.end(); it !=
e;
+ ++it)
+ {
+ i += (*it)->Marshal(k, i, nVectorWidth, pProgram);
}
}
/// Handling a Double Vector that is used as a sliding window input
/// Performs parallel reduction based on given operator
template<class Base>
-class ParallelReductionVectorRef: public Base
+class ParallelReductionVectorRef : public Base
{
public:
- ParallelReductionVectorRef(const std::string &s,
- FormulaTreeNodeRef ft, boost::shared_ptr<SlidingFunctionBase> &CodeGen,
- int index=0):
+ ParallelReductionVectorRef( const std::string& s,
+ FormulaTreeNodeRef ft, boost::shared_ptr<SlidingFunctionBase>& CodeGen,
+ int index = 0 ) :
Base(s, ft, index), mpCodeGen(CodeGen), mpClmem2(NULL)
{
- FormulaToken *t = ft->GetFormulaToken();
+ FormulaToken* t = ft->GetFormulaToken();
if (t->GetType() != formula::svDoubleVectorRef)
throw Unhandled();
- mpDVR = static_cast<const formula::DoubleVectorRefToken *>(t);
+ mpDVR = static_cast<const formula::DoubleVectorRefToken*>(t);
bIsStartFixed = mpDVR->IsStartFixed();
bIsEndFixed = mpDVR->IsEndFixed();
}
/// Emit the definition for the auxiliary reduction kernel
- virtual void GenSlidingWindowFunction(std::stringstream &ss) {
- if ( !dynamic_cast<OpAverage*>(mpCodeGen.get()))
- {
- std::string name = Base::GetName();
- ss << "__kernel void "<<name;
- ss << "_reduction(__global double* A, "
- "__global double *result,int arrayLength,int windowSize){\n";
- ss << " double tmp, current_result =" <<
- mpCodeGen->GetBottom();
- ss << ";\n";
- ss << " int writePos = get_group_id(1);\n";
- ss << " int lidx = get_local_id(0);\n";
- ss << " __local double shm_buf[256];\n";
- if (mpDVR->IsStartFixed())
- ss << " int offset = 0;\n";
- else // if (!mpDVR->IsStartFixed())
- ss << " int offset = get_group_id(1);\n";
- if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
- ss << " int end = windowSize;\n";
- else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
- ss << " int end = offset + windowSize;\n";
- else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
- ss << " int end = windowSize + get_group_id(1);\n";
- else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
- ss << " int end = windowSize;\n";
- ss << " end = min(end, arrayLength);\n";
-
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " int loop = arrayLength/512 + 1;\n";
- ss << " for (int l=0; l<loop; l++){\n";
- ss << " tmp = "<< mpCodeGen->GetBottom() << ";\n";
- ss << " int loopOffset = l*512;\n";
- ss << " if((loopOffset + lidx + offset + 256) < end) {\n";
- ss << " tmp = legalize(" << mpCodeGen->Gen2(
- "A[loopOffset + lidx + offset]", "tmp") <<", tmp);\n";
- ss << " tmp = legalize(" << mpCodeGen->Gen2(
- "A[loopOffset + lidx + offset + 256]", "tmp") <<", tmp);\n";
- ss << " } else if ((loopOffset + lidx + offset) < end)\n";
- ss << " tmp = legalize(" << mpCodeGen->Gen2(
- "A[loopOffset + lidx + offset]", "tmp") <<", tmp);\n";
- ss << " shm_buf[lidx] = tmp;\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " for (int i = 128; i >0; i/=2) {\n";
- ss << " if (lidx < i)\n";
- ss << " shm_buf[lidx] = ";
- // Special case count
- if (dynamic_cast<OpCount*>(mpCodeGen.get()))
- ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
+ virtual void GenSlidingWindowFunction( std::stringstream& ss )
+ {
+ if (!dynamic_cast<OpAverage*>(mpCodeGen.get()))
+ {
+ std::string name = Base::GetName();
+ ss << "__kernel void " << name;
+ ss << "_reduction(__global double* A, "
+ "__global double *result,int arrayLength,int windowSize){\n";
+ ss << " double tmp, current_result =" <<
+ mpCodeGen->GetBottom();
+ ss << ";\n";
+ ss << " int writePos = get_group_id(1);\n";
+ ss << " int lidx = get_local_id(0);\n";
+ ss << " __local double shm_buf[256];\n";
+ if (mpDVR->IsStartFixed())
+ ss << " int offset = 0;\n";
+ else // if (!mpDVR->IsStartFixed())
+ ss << " int offset = get_group_id(1);\n";
+ if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
+ ss << " int end = windowSize;\n";
+ else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ ss << " int end = offset + windowSize;\n";
+ else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ ss << " int end = windowSize + get_group_id(1);\n";
+ else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
+ ss << " int end = windowSize;\n";
+ ss << " end = min(end, arrayLength);\n";
+
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " int loop = arrayLength/512 + 1;\n";
+ ss << " for (int l=0; l<loop; l++){\n";
+ ss << " tmp = " << mpCodeGen->GetBottom() << ";\n";
+ ss << " int loopOffset = l*512;\n";
+ ss << " if((loopOffset + lidx + offset + 256) < end) {\n";
+ ss << " tmp = legalize(" << mpCodeGen->Gen2(
+ "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n";
+ ss << " tmp = legalize(" << mpCodeGen->Gen2(
+ "A[loopOffset + lidx + offset + 256]", "tmp") << ", tmp);\n";
+ ss << " } else if ((loopOffset + lidx + offset) < end)\n";
+ ss << " tmp = legalize(" << mpCodeGen->Gen2(
+ "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n";
+ ss << " shm_buf[lidx] = tmp;\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " for (int i = 128; i >0; i/=2) {\n";
+ ss << " if (lidx < i)\n";
+ ss << " shm_buf[lidx] = ";
+ // Special case count
+ if (dynamic_cast<OpCount*>(mpCodeGen.get()))
+ ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
+ else
+ ss << mpCodeGen->Gen2("shm_buf[lidx]", "shm_buf[lidx + i]") <<
";\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " }\n";
+ ss << " if (lidx == 0)\n";
+ ss << " current_result =";
+ if (dynamic_cast<OpCount*>(mpCodeGen.get()))
+ ss << "current_result + shm_buf[0]";
+ else
+ ss << mpCodeGen->Gen2("current_result", "shm_buf[0]");
+ ss << ";\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " }\n";
+ ss << " if (lidx == 0)\n";
+ ss << " result[writePos] = current_result;\n";
+ ss << "}\n";
+ }
else
- ss << mpCodeGen->Gen2("shm_buf[lidx]", "shm_buf[lidx + i]")<<";\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " }\n";
- ss << " if (lidx == 0)\n";
- ss << " current_result =";
- if (dynamic_cast<OpCount*>(mpCodeGen.get()))
+ {
+ std::string name = Base::GetName();
+ /*sum reduction*/
+ ss << "__kernel void " << name << "_sum";
+ ss << "_reduction(__global double* A, "
+ "__global double *result,int arrayLength,int windowSize){\n";
+ ss << " double tmp, current_result =" <<
+ mpCodeGen->GetBottom();
+ ss << ";\n";
+ ss << " int writePos = get_group_id(1);\n";
+ ss << " int lidx = get_local_id(0);\n";
+ ss << " __local double shm_buf[256];\n";
+ if (mpDVR->IsStartFixed())
+ ss << " int offset = 0;\n";
+ else // if (!mpDVR->IsStartFixed())
+ ss << " int offset = get_group_id(1);\n";
+ if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
+ ss << " int end = windowSize;\n";
+ else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ ss << " int end = offset + windowSize;\n";
+ else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ ss << " int end = windowSize + get_group_id(1);\n";
+ else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
+ ss << " int end = windowSize;\n";
+ ss << " end = min(end, arrayLength);\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " int loop = arrayLength/512 + 1;\n";
+ ss << " for (int l=0; l<loop; l++){\n";
+ ss << " tmp = " << mpCodeGen->GetBottom() << ";\n";
+ ss << " int loopOffset = l*512;\n";
+ ss << " if((loopOffset + lidx + offset + 256) < end) {\n";
+ ss << " tmp = legalize(";
+ ss << "(A[loopOffset + lidx + offset]+ tmp)";
+ ss << ", tmp);\n";
+ ss << " tmp = legalize((A[loopOffset + lidx + offset +
256]+ tmp)";
+ ss << ", tmp);\n";
+ ss << " } else if ((loopOffset + lidx + offset) < end)\n";
+ ss << " tmp = legalize((A[loopOffset + lidx + offset] +
tmp)";
+ ss << ", tmp);\n";
+ ss << " shm_buf[lidx] = tmp;\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " for (int i = 128; i >0; i/=2) {\n";
+ ss << " if (lidx < i)\n";
+ ss << " shm_buf[lidx] = ";
+ ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " }\n";
+ ss << " if (lidx == 0)\n";
+ ss << " current_result =";
ss << "current_result + shm_buf[0]";
- else
- ss << mpCodeGen->Gen2("current_result", "shm_buf[0]");
- ss << ";\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " }\n";
- ss << " if (lidx == 0)\n";
- ss << " result[writePos] = current_result;\n";
- ss << "}\n";
- }
- else{
- std::string name = Base::GetName();
- /*sum reduction*/
- ss << "__kernel void "<<name<<"_sum";
- ss << "_reduction(__global double* A, "
- "__global double *result,int arrayLength,int windowSize){\n";
- ss << " double tmp, current_result =" <<
- mpCodeGen->GetBottom();
- ss << ";\n";
- ss << " int writePos = get_group_id(1);\n";
- ss << " int lidx = get_local_id(0);\n";
- ss << " __local double shm_buf[256];\n";
- if (mpDVR->IsStartFixed())
- ss << " int offset = 0;\n";
- else // if (!mpDVR->IsStartFixed())
- ss << " int offset = get_group_id(1);\n";
- if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
- ss << " int end = windowSize;\n";
- else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
- ss << " int end = offset + windowSize;\n";
- else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
- ss << " int end = windowSize + get_group_id(1);\n";
- else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
- ss << " int end = windowSize;\n";
- ss << " end = min(end, arrayLength);\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " int loop = arrayLength/512 + 1;\n";
- ss << " for (int l=0; l<loop; l++){\n";
- ss << " tmp = "<< mpCodeGen->GetBottom() << ";\n";
- ss << " int loopOffset = l*512;\n";
- ss << " if((loopOffset + lidx + offset + 256) < end) {\n";
- ss << " tmp = legalize(";
- ss << "(A[loopOffset + lidx + offset]+ tmp)";
- ss << ", tmp);\n";
- ss << " tmp = legalize((A[loopOffset + lidx + offset + 256]+
tmp)";
- ss << ", tmp);\n";
- ss << " } else if ((loopOffset + lidx + offset) < end)\n";
- ss << " tmp = legalize((A[loopOffset + lidx + offset] + tmp)";
- ss << ", tmp);\n";
- ss << " shm_buf[lidx] = tmp;\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " for (int i = 128; i >0; i/=2) {\n";
- ss << " if (lidx < i)\n";
- ss << " shm_buf[lidx] = ";
- ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " }\n";
- ss << " if (lidx == 0)\n";
- ss << " current_result =";
- ss << "current_result + shm_buf[0]";
- ss << ";\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " }\n";
- ss << " if (lidx == 0)\n";
- ss << " result[writePos] = current_result;\n";
- ss << "}\n";
- /*count reduction*/
- ss << "__kernel void "<<name<<"_count";
- ss << "_reduction(__global double* A, "
- "__global double *result,int arrayLength,int windowSize){\n";
- ss << " double tmp, current_result =" <<
- mpCodeGen->GetBottom();
- ss << ";\n";
- ss << " int writePos = get_group_id(1);\n";
- ss << " int lidx = get_local_id(0);\n";
- ss << " __local double shm_buf[256];\n";
- if (mpDVR->IsStartFixed())
- ss << " int offset = 0;\n";
- else // if (!mpDVR->IsStartFixed())
- ss << " int offset = get_group_id(1);\n";
- if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
- ss << " int end = windowSize;\n";
- else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
- ss << " int end = offset + windowSize;\n";
- else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
- ss << " int end = windowSize + get_group_id(1);\n";
- else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
- ss << " int end = windowSize;\n";
- ss << " end = min(end, arrayLength);\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " int loop = arrayLength/512 + 1;\n";
- ss << " for (int l=0; l<loop; l++){\n";
- ss << " tmp = "<< mpCodeGen->GetBottom() << ";\n";
- ss << " int loopOffset = l*512;\n";
- ss << " if((loopOffset + lidx + offset + 256) < end) {\n";
- ss << " tmp = legalize((isNan(A[loopOffset + lidx +
offset])?tmp:tmp+1.0)";
- ss << ", tmp);\n";
- ss << " tmp = legalize((isNan(A[loopOffset + lidx +
offset+256])?tmp:tmp+1.0)";
- ss << ", tmp);\n";
- ss << " } else if ((loopOffset + lidx + offset) < end)\n";
- ss << " tmp = legalize((isNan(A[loopOffset + lidx +
offset])?tmp:tmp+1.0)";
- ss << ", tmp);\n";
- ss << " shm_buf[lidx] = tmp;\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " for (int i = 128; i >0; i/=2) {\n";
- ss << " if (lidx < i)\n";
- ss << " shm_buf[lidx] = ";
- ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " }\n";
- ss << " if (lidx == 0)\n";
- ss << " current_result =";
- ss << "current_result + shm_buf[0];";
- ss << ";\n";
- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
- ss << " }\n";
- ss << " if (lidx == 0)\n";
- ss << " result[writePos] = current_result;\n";
- ss << "}\n";
- }
+ ss << ";\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " }\n";
+ ss << " if (lidx == 0)\n";
+ ss << " result[writePos] = current_result;\n";
+ ss << "}\n";
+ /*count reduction*/
+ ss << "__kernel void " << name << "_count";
+ ss << "_reduction(__global double* A, "
+ "__global double *result,int arrayLength,int windowSize){\n";
+ ss << " double tmp, current_result =" <<
+ mpCodeGen->GetBottom();
+ ss << ";\n";
+ ss << " int writePos = get_group_id(1);\n";
+ ss << " int lidx = get_local_id(0);\n";
+ ss << " __local double shm_buf[256];\n";
+ if (mpDVR->IsStartFixed())
+ ss << " int offset = 0;\n";
+ else // if (!mpDVR->IsStartFixed())
+ ss << " int offset = get_group_id(1);\n";
+ if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
+ ss << " int end = windowSize;\n";
+ else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ ss << " int end = offset + windowSize;\n";
+ else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
+ ss << " int end = windowSize + get_group_id(1);\n";
+ else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
+ ss << " int end = windowSize;\n";
+ ss << " end = min(end, arrayLength);\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " int loop = arrayLength/512 + 1;\n";
+ ss << " for (int l=0; l<loop; l++){\n";
+ ss << " tmp = " << mpCodeGen->GetBottom() << ";\n";
+ ss << " int loopOffset = l*512;\n";
+ ss << " if((loopOffset + lidx + offset + 256) < end) {\n";
+ ss << " tmp = legalize((isNan(A[loopOffset + lidx +
offset])?tmp:tmp+1.0)";
+ ss << ", tmp);\n";
+ ss << " tmp = legalize((isNan(A[loopOffset + lidx +
offset+256])?tmp:tmp+1.0)";
+ ss << ", tmp);\n";
+ ss << " } else if ((loopOffset + lidx + offset) < end)\n";
+ ss << " tmp = legalize((isNan(A[loopOffset + lidx +
offset])?tmp:tmp+1.0)";
+ ss << ", tmp);\n";
+ ss << " shm_buf[lidx] = tmp;\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " for (int i = 128; i >0; i/=2) {\n";
+ ss << " if (lidx < i)\n";
+ ss << " shm_buf[lidx] = ";
+ ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " }\n";
+ ss << " if (lidx == 0)\n";
+ ss << " current_result =";
+ ss << "current_result + shm_buf[0];";
+ ss << ";\n";
+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
+ ss << " }\n";
+ ss << " if (lidx == 0)\n";
+ ss << " result[writePos] = current_result;\n";
+ ss << "}\n";
+ }
}
- virtual std::string GenSlidingWindowDeclRef(bool=false) const
+ virtual std::string GenSlidingWindowDeclRef( bool = false ) const
{
std::stringstream ss;
if (!bIsStartFixed && !bIsEndFixed)
@@ -1047,19 +1085,19 @@ public:
}
/// Controls how the elements in the DoubleVectorRef are traversed
virtual size_t GenReductionLoopHeader(
- std::stringstream &ss, bool &needBody)
+ std::stringstream& ss, bool& needBody )
{
assert(mpDVR);
size_t nCurWindowSize = mpDVR->GetRefRowSize();
std::string temp = Base::GetName() + "[gid0]";
ss << "tmp = ";
// Special case count
- if ( dynamic_cast<OpAverage*>(mpCodeGen.get()))
+ if (dynamic_cast<OpAverage*>(mpCodeGen.get()))
{
- ss << mpCodeGen->Gen2(temp, "tmp")<<";\n";
- ss <<"nCount = nCount-1;\n";
- ss <<"nCount = nCount +";/*re-assign nCount from count reduction*/
- ss << Base::GetName()<<"[gid0+"<<SymbolTable::nR<<"]"<<";\n";
+ ss << mpCodeGen->Gen2(temp, "tmp") << ";\n";
+ ss << "nCount = nCount-1;\n";
+ ss << "nCount = nCount +"; /*re-assign nCount from count
reduction*/
+ ss << Base::GetName() << "[gid0+" << SymbolTable::nR << "]" <<
";\n";
}
else if (dynamic_cast<OpCount*>(mpCodeGen.get()))
ss << temp << "+ tmp";
@@ -1070,7 +1108,7 @@ public:
return nCurWindowSize;
}
- virtual size_t Marshal(cl_kernel k, int argno, int w, cl_program mpProgram)
+ virtual size_t Marshal( cl_kernel k, int argno, int w, cl_program
mpProgram )
{
assert(Base::mpClmem == NULL);
// Obtain cl context
@@ -1082,24 +1120,24 @@ public:
// create clmem buffer
if (mpDVR->GetArrays()[Base::mnIndex].mpNumericArray == NULL)
throw Unhandled();
- double *pHostBuffer = const_cast<double*>(
- mpDVR->GetArrays()[Base::mnIndex].mpNumericArray);
+ double* pHostBuffer = const_cast<double*>(
+ mpDVR->GetArrays()[Base::mnIndex].mpNumericArray);
size_t szHostBuffer = nInput * sizeof(double);
Base::mpClmem = clCreateBuffer(kEnv.mpkContext,
- (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR,
- szHostBuffer,
- pHostBuffer, &err);
+ (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+ szHostBuffer,
+ pHostBuffer, &err);
mpClmem2 = clCreateBuffer(kEnv.mpkContext,
- CL_MEM_READ_WRITE|CL_MEM_ALLOC_HOST_PTR,
- sizeof(double)*w, NULL, NULL);
+ CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
+ sizeof(double) * w, NULL, NULL);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
// reproduce the reduction function name
std::string kernelName;
- if ( !dynamic_cast<OpAverage*>(mpCodeGen.get()))
- kernelName = Base::GetName() + "_reduction";
+ if (!dynamic_cast<OpAverage*>(mpCodeGen.get()))
+ kernelName = Base::GetName() + "_reduction";
else
- kernelName = Base::GetName() + "_sum_reduction";
+ kernelName = Base::GetName() + "_sum_reduction";
cl_kernel redKernel = clCreateKernel(mpProgram, kernelName.c_str(),
&err);
if (err != CL_SUCCESS)
throw OpenCLError(err, __FILE__, __LINE__);
@@ -1107,11 +1145,11 @@ public:
// TODO(Wei Wei): use unique name for kernel
cl_mem buf = Base::GetCLBuffer();
err = clSetKernelArg(redKernel, 0, sizeof(cl_mem),
- (void *)&buf);
+ (void*)&buf);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
- err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), (void *)&mpClmem2);
+ err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), (void*)&mpClmem2);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
@@ -1124,28 +1162,28 @@ public:
throw OpenCLError(err, __FILE__, __LINE__);
// set work group size and execute
- size_t global_work_size[] = {256, (size_t)w };
- size_t local_work_size[] = {256, 1};
+ size_t global_work_size[] = { 256, (size_t)w };
+ size_t local_work_size[] = { 256, 1 };
err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL,
- global_work_size, local_work_size, 0, NULL, NULL);
+ global_work_size, local_work_size, 0, NULL, NULL);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
err = clFinish(kEnv.mpkCmdQueue);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
- if ( dynamic_cast<OpAverage*>(mpCodeGen.get()))
+ if (dynamic_cast<OpAverage*>(mpCodeGen.get()))
{
- /*average need more reduction kernel for count computing*/
- boost::scoped_array<double> pAllBuffer(new double[2*w]);
- double *resbuf = (double*)clEnqueueMapBuffer(kEnv.mpkCmdQueue,
- mpClmem2,
- CL_TRUE, CL_MAP_READ, 0,
- sizeof(double)*w, 0, NULL, NULL,
- &err);
+ /*average need more reduction kernel for count computing*/
+ boost::scoped_array<double> pAllBuffer(new double[2 * w]);
+ double* resbuf = (double*)clEnqueueMapBuffer(kEnv.mpkCmdQueue,
+ mpClmem2,
+ CL_TRUE, CL_MAP_READ, 0,
+ sizeof(double) * w, 0, NULL, NULL,
+ &err);
if (err != CL_SUCCESS)
throw OpenCLError(err, __FILE__, __LINE__);
- for (int i=0 ; i < w; i++)
+ for (int i = 0; i < w; i++)
pAllBuffer[i] = resbuf[i];
err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf,
0, NULL, NULL);
if (err != CL_SUCCESS)
@@ -1158,11 +1196,11 @@ public:
// set kernel arg of reduction kernel
buf = Base::GetCLBuffer();
err = clSetKernelArg(redKernel, 0, sizeof(cl_mem),
- (void *)&buf);
+ (void*)&buf);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
- err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), (void
*)&mpClmem2);
+ err = clSetKernelArg(redKernel, 1, sizeof(cl_mem),
(void*)&mpClmem2);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
@@ -1175,10 +1213,10 @@ public:
throw OpenCLError(err, __FILE__, __LINE__);
// set work group size and execute
- size_t global_work_size1[] = {256, (size_t)w };
- size_t local_work_size1[] = {256, 1};
+ size_t global_work_size1[] = { 256, (size_t)w };
+ size_t local_work_size1[] = { 256, 1 };
err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL,
- global_work_size1, local_work_size1, 0, NULL, NULL);
+ global_work_size1, local_work_size1, 0, NULL, NULL);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
err = clFinish(kEnv.mpkCmdQueue);
@@ -1187,12 +1225,12 @@ public:
resbuf = (double*)clEnqueueMapBuffer(kEnv.mpkCmdQueue,
mpClmem2,
CL_TRUE, CL_MAP_READ, 0,
- sizeof(double)*w, 0, NULL, NULL,
+ sizeof(double) * w, 0, NULL, NULL,
&err);
if (err != CL_SUCCESS)
throw OpenCLError(err, __FILE__, __LINE__);
- for (int i=0 ; i < w; i++)
- pAllBuffer[i+w] = resbuf[i];
+ for (int i = 0; i < w; i++)
+ pAllBuffer[i + w] = resbuf[i];
err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf,
0, NULL, NULL);
if (mpClmem2)
{
@@ -1200,8 +1238,8 @@ public:
mpClmem2 = NULL;
}
mpClmem2 = clCreateBuffer(kEnv.mpkContext,
- (cl_mem_flags) CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
- w*sizeof(double)*2, pAllBuffer.get(), &err);
+ (cl_mem_flags)CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+ w * sizeof(double) * 2, pAllBuffer.get(), &err);
if (CL_SUCCESS != err)
throw OpenCLError(err, __FILE__, __LINE__);
}
@@ -1220,35 +1258,35 @@ public:
}
}
- size_t GetArrayLength(void) const {return mpDVR->GetArrayLength(); }
+ size_t GetArrayLength() const { return mpDVR->GetArrayLength(); }
- size_t GetWindowSize(void) const {return mpDVR->GetRefRowSize(); }
+ size_t GetWindowSize() const { return mpDVR->GetRefRowSize(); }
- size_t GetStartFixed(void) const {return bIsStartFixed; }
+ size_t GetStartFixed() const { return bIsStartFixed; }
- size_t GetEndFixed(void) const {return bIsEndFixed; }
+ size_t GetEndFixed() const { return bIsEndFixed; }
protected:
bool bIsStartFixed, bIsEndFixed;
- const formula::DoubleVectorRefToken *mpDVR;
+ const formula::DoubleVectorRefToken* mpDVR;
// from parent nodes
boost::shared_ptr<SlidingFunctionBase> mpCodeGen;
// controls whether to invoke the reduction kernel during marshaling or not
cl_mem mpClmem2;
};
-class Reduction: public SlidingFunctionBase
+class Reduction : public SlidingFunctionBase
{
public:
typedef DynamicKernelSlidingArgument<VectorRef> NumericRange;
typedef DynamicKernelSlidingArgument<DynamicKernelStringArgument>
StringRange;
typedef ParallelReductionVectorRef<VectorRef> ParallelNumericRange;
- virtual void GenSlidingWindowFunction(std::stringstream &ss,
- const std::string &sSymName, SubArguments &vSubArguments)
SAL_OVERRIDE
+ virtual void GenSlidingWindowFunction( std::stringstream& ss,
+ const std::string& sSymName, SubArguments& vSubArguments ) SAL_OVERRIDE
{
ss << "\ndouble " << sSymName;
- ss << "_"<< BinFuncName() <<"(";
+ ss << "_" << BinFuncName() << "(";
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
if (i)
@@ -1256,7 +1294,7 @@ public:
vSubArguments[i]->GenSlidingWindowDecl(ss);
}
ss << ") {\n";
- ss << "double tmp = " << GetBottom() <<";\n";
+ ss << "double tmp = " << GetBottom() << ";\n";
ss << "int gid0 = get_global_id(0);\n";
if (isAverage())
ss << "int nCount = 0;\n";
@@ -1264,35 +1302,38 @@ public:
unsigned i = vSubArguments.size();
while (i--)
{
- if (NumericRange *NR =
- dynamic_cast<NumericRange *> (vSubArguments[i].get()))
+ if (NumericRange* NR =
+ dynamic_cast<NumericRange*>(vSubArguments[i].get()))
{
- bool needBody;NR->GenReductionLoopHeader(ss, needBody);if
(needBody == false) continue;
+ bool needBody; NR->GenReductionLoopHeader(ss, needBody); if
(needBody == false)
+ continue;
}
- else if (ParallelNumericRange *PNR =
- dynamic_cast<ParallelNumericRange *>
(vSubArguments[i].get()))
+ else if (ParallelNumericRange* PNR =
+ dynamic_cast<ParallelNumericRange*>(vSubArguments[i].get()))
{
//did not handle yet
- bool needBody;PNR->GenReductionLoopHeader(ss, needBody);if
(needBody == false) continue;
+ bool needBody; PNR->GenReductionLoopHeader(ss, needBody); if
(needBody == false)
+ continue;
}
- else if (StringRange *SR =
- dynamic_cast<StringRange *> (vSubArguments[i].get()))
+ else if (StringRange* SR =
+ dynamic_cast<StringRange*>(vSubArguments[i].get()))
{
//did not handle yet
bool needBody;
SR->GenReductionLoopHeader(ss, needBody);
- if (needBody == false) continue;
+ if (needBody == false)
+ continue;
}
else
{
- FormulaToken *pCur = vSubArguments[i]->GetFormulaToken();
+ FormulaToken* pCur = vSubArguments[i]->GetFormulaToken();
assert(pCur);
assert(pCur->GetType() != formula::svDoubleVectorRef);
if (pCur->GetType() == formula::svSingleVectorRef)
{
const formula::SingleVectorRefToken* pSVR =
- static_cast< const formula::SingleVectorRefToken*
>(pCur);
+ static_cast<const
formula::SingleVectorRefToken*>(pCur);
ss << "if (gid0 < " << pSVR->GetArrayLength() << "){\n";
}
else if (pCur->GetType() == formula::svDouble)
@@ -1300,13 +1341,13 @@ public:
ss << "{\n";
}
}
- if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
+ if (ocPush == vSubArguments[i]->GetFormulaToken()->GetOpCode())
{
ss << "tmpBottom = " << GetBottom() << ";\n";
ss << "if (isNan(";
ss << vSubArguments[i]->GenSlidingWindowDeclRef();
ss << "))\n";
- if ( ZeroReturnZero() )
+ if (ZeroReturnZero())
ss << " return 0;\n";
else
{
@@ -1319,8 +1360,8 @@ public:
ss << ";\n";
ss << " }\n";
ss << "}\n";
- if ( vSubArguments[i]->GetFormulaToken()->GetType() ==
- formula::svSingleVectorRef&& ZeroReturnZero() )
+ if (vSubArguments[i]->GetFormulaToken()->GetType() ==
+ formula::svSingleVectorRef && ZeroReturnZero())
{
ss << "else{\n";
ss << " return 0;\n";
@@ -1345,14 +1386,14 @@ public:
};
// Strictly binary operators
-class Binary: public SlidingFunctionBase
+class Binary : public SlidingFunctionBase
{
public:
- virtual void GenSlidingWindowFunction(std::stringstream &ss,
- const std::string &sSymName, SubArguments &vSubArguments)
SAL_OVERRIDE
+ virtual void GenSlidingWindowFunction( std::stringstream& ss,
+ const std::string& sSymName, SubArguments& vSubArguments ) SAL_OVERRIDE
{
ss << "\ndouble " << sSymName;
- ss << "_"<< BinFuncName() <<"(";
+ ss << "_" << BinFuncName() << "(";
assert(vSubArguments.size() == 2);
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
@@ -1364,41 +1405,41 @@ public:
ss << "int gid0 = get_global_id(0), i = 0;\n\t";
ss << "double tmp = ";
ss << Gen2(vSubArguments[0]->GenSlidingWindowDeclRef(false),
- vSubArguments[1]->GenSlidingWindowDeclRef(false)) << ";\n\t";
+ vSubArguments[1]->GenSlidingWindowDeclRef(false)) << ";\n\t";
ss << "return tmp;\n}";
}
virtual bool takeString() const SAL_OVERRIDE { return true; }
virtual bool takeNumeric() const SAL_OVERRIDE { return true; }
};
-class SumOfProduct: public SlidingFunctionBase
+class SumOfProduct : public SlidingFunctionBase
{
public:
- virtual void GenSlidingWindowFunction(std::stringstream &ss,
- const std::string &sSymName, SubArguments &vSubArguments)
SAL_OVERRIDE
+ virtual void GenSlidingWindowFunction( std::stringstream& ss,
+ const std::string& sSymName, SubArguments& vSubArguments ) SAL_OVERRIDE
{
size_t nCurWindowSize = 0;
- FormulaToken *tmpCur = NULL;
- const formula::DoubleVectorRefToken *pCurDVR = NULL;
+ FormulaToken* tmpCur = NULL;
+ const formula::DoubleVectorRefToken* pCurDVR = NULL;
ss << "\ndouble " << sSymName;
- ss << "_"<< BinFuncName() <<"(";
+ ss << "_" << BinFuncName() << "(";
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
if (i)
ss << ",";
vSubArguments[i]->GenSlidingWindowDecl(ss);
size_t nCurChildWindowSize = vSubArguments[i]->GetWindowSize();
- nCurWindowSize = (nCurWindowSize < nCurChildWindowSize)?
- nCurChildWindowSize:nCurWindowSize;
+ nCurWindowSize = (nCurWindowSize < nCurChildWindowSize) ?
+ nCurChildWindowSize : nCurWindowSize;
tmpCur = vSubArguments[i]->GetFormulaToken();
- if ( ocPush==tmpCur->GetOpCode() )
+ if (ocPush == tmpCur->GetOpCode())
{
pCurDVR = static_cast<
- const formula::DoubleVectorRefToken*>(tmpCur);
- if ( !
- ( (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
- || (pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) )
+ const
formula::DoubleVectorRefToken*>(tmpCur);
+ if (!
+ ((!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
+ || (pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()))
)
throw Unhandled();
}
@@ -1408,20 +1449,20 @@ public:
ss << " int gid0 = get_global_id(0);\n";
#ifndef UNROLLING_FACTOR
ss << " int i ;\n";
- ss << " for (i = 0; i < "<< nCurWindowSize <<"; i++)\n";
+ ss << " for (i = 0; i < " << nCurWindowSize << "; i++)\n";
ss << " {\n";
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
tmpCur = vSubArguments[i]->GetFormulaToken();
- if(ocPush==tmpCur->GetOpCode())
+ if (ocPush == tmpCur->GetOpCode())
{
- pCurDVR= static_cast<
- const formula::DoubleVectorRefToken *>(tmpCur);
- if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
+ pCurDVR = static_cast<
+ const
formula::DoubleVectorRefToken*>(tmpCur);
+ if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
{
- ss << " int currentCount";
- ss << i;
- ss <<" =i+gid0+1;\n";
+ ss << " int currentCount";
+ ss << i;
+ ss << " =i+gid0+1;\n";
}
else
{
@@ -1429,7 +1470,7 @@ public:
ss << i;
ss << " =i+1;\n";
}
- }
+ }
}
ss << " tmp += fsum(";
for (unsigned i = 0; i < vSubArguments.size(); i++)
@@ -1437,29 +1478,29 @@ public:
if (i)
ss << "*";
#ifdef ISNAN
- if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
+ if (ocPush == vSubArguments[i]->GetFormulaToken()->GetOpCode())
{
- ss <<"(";
- ss <<"(currentCount";
+ ss << "(";
+ ss << "(currentCount";
ss << i;
- ss<< ">";
- if(vSubArguments[i]->GetFormulaToken()->GetType() ==
- formula::svSingleVectorRef)
+ ss << ">";
+ if (vSubArguments[i]->GetFormulaToken()->GetType() ==
+ formula::svSingleVectorRef)
{
const formula::SingleVectorRefToken* pSVR =
- static_cast< const formula::SingleVectorRefToken*>
- (vSubArguments[i]->GetFormulaToken());
- ss<<pSVR->GetArrayLength();
+ static_cast<const formula::SingleVectorRefToken*>
+ (vSubArguments[i]->GetFormulaToken());
+ ss << pSVR->GetArrayLength();
}
- else if(vSubArguments[i]->GetFormulaToken()->GetType() ==
- formula::svDoubleVectorRef)
+ else if (vSubArguments[i]->GetFormulaToken()->GetType() ==
+ formula::svDoubleVectorRef)
{
const formula::DoubleVectorRefToken* pSVR =
- static_cast< const formula::DoubleVectorRefToken*>
- (vSubArguments[i]->GetFormulaToken());
- ss<<pSVR->GetArrayLength();
+ static_cast<const formula::DoubleVectorRefToken*>
+ (vSubArguments[i]->GetFormulaToken());
+ ss << pSVR->GetArrayLength();
}
- ss << ")||isNan("<<vSubArguments[i]
+ ss << ")||isNan(" << vSubArguments[i]
->GenSlidingWindowDeclRef(true);
ss << ")?0:";
ss << vSubArguments[i]->GenSlidingWindowDeclRef(true);
@@ -1479,28 +1520,31 @@ public:
#ifdef UNROLLING_FACTOR
ss << "\tint i;\n\t";
ss << "int currentCount0;\n";
- for ( unsigned i = 0; i < vSubArguments.size()-1; i++)
- ss << "int currentCount"<<i+1<<";\n";
- std::stringstream temp3,temp4;
+ for (unsigned i = 0; i < vSubArguments.size() - 1; i++)
+ ss << "int currentCount" << i + 1 << ";\n";
+ std::stringstream temp3, temp4;
int outLoopSize = UNROLLING_FACTOR;
- if (nCurWindowSize/outLoopSize != 0){
+ if (nCurWindowSize / outLoopSize != 0)
+ {
ss << "for(int outLoop=0; outLoop<" <<
- nCurWindowSize/outLoopSize<< "; outLoop++){\n\t";
- for(int count=0; count < outLoopSize; count++){
- ss << "i = outLoop*"<<outLoopSize<<"+"<<count<<";\n";
- if(count==0){
+ nCurWindowSize / outLoopSize << "; outLoop++){\n\t";
+ for (int count = 0; count < outLoopSize; count++)
+ {
+ ss << "i = outLoop*" << outLoopSize << "+" << count << ";\n";
+ if (count == 0)
+ {
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
tmpCur = vSubArguments[i]->GetFormulaToken();
- if(ocPush==tmpCur->GetOpCode())
+ if (ocPush == tmpCur->GetOpCode())
{
- pCurDVR= static_cast<
- const formula::DoubleVectorRefToken *>(tmpCur);
- if(!pCurDVR->IsStartFixed() &&
!pCurDVR->IsEndFixed())
+ pCurDVR = static_cast<
+ const
formula::DoubleVectorRefToken*>(tmpCur);
+ if (!pCurDVR->IsStartFixed() &&
!pCurDVR->IsEndFixed())
{
temp3 << " currentCount";
temp3 << i;
- temp3 <<" =i+gid0+1;\n";
+ temp3 << " =i+gid0+1;\n";
}
else
{
@@ -1512,37 +1556,41 @@ public:
}
temp3 << "tmp = fsum(";
- for (unsigned i = 0; i < vSubArguments.size(); i++){
+ for (unsigned i = 0; i < vSubArguments.size(); i++)
+ {
if (i)
temp3 << "*";
-
if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode()){
- temp3 <<"(";
- temp3 <<"(currentCount";
+ if (ocPush ==
vSubArguments[i]->GetFormulaToken()->GetOpCode())
+ {
+ temp3 << "(";
+ temp3 << "(currentCount";
temp3 << i;
temp3 << ">";
- if(vSubArguments[i]->GetFormulaToken()->GetType()
==
- formula::svSingleVectorRef){
+ if (vSubArguments[i]->GetFormulaToken()->GetType()
==
+ formula::svSingleVectorRef)
+ {
const formula::SingleVectorRefToken* pSVR =
- static_cast< const
formula::SingleVectorRefToken*>
+ static_cast<const
formula::SingleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken());
- temp3<<pSVR->GetArrayLength();
- temp3 << ")||isNan("<<vSubArguments[i]
- ->GenSlidingWindowDeclRef();
- temp3 << ")?0:";
- temp3 <<
vSubArguments[i]->GenSlidingWindowDeclRef();
- temp3 << ")";
+ temp3 << pSVR->GetArrayLength();
+ temp3 << ")||isNan(" << vSubArguments[i]
+ ->GenSlidingWindowDeclRef();
+ temp3 << ")?0:";
+ temp3 <<
vSubArguments[i]->GenSlidingWindowDeclRef();
+ temp3 << ")";
}
- else
if(vSubArguments[i]->GetFormulaToken()->GetType() ==
- formula::svDoubleVectorRef){
+ else if
(vSubArguments[i]->GetFormulaToken()->GetType() ==
+ formula::svDoubleVectorRef)
+ {
const formula::DoubleVectorRefToken* pSVR =
- static_cast< const
formula::DoubleVectorRefToken*>
+ static_cast<const
formula::DoubleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken());
- temp3<<pSVR->GetArrayLength();
- temp3 << ")||isNan("<<vSubArguments[i]
- ->GenSlidingWindowDeclRef(true);
- temp3 << ")?0:";
- temp3 <<
vSubArguments[i]->GenSlidingWindowDeclRef(true);
- temp3 << ")";
+ temp3 << pSVR->GetArrayLength();
+ temp3 << ")||isNan(" << vSubArguments[i]
+ ->GenSlidingWindowDeclRef(true);
+ temp3 << ")?0:";
+ temp3 <<
vSubArguments[i]->GenSlidingWindowDeclRef(true);
+ temp3 << ")";
}
}
@@ -1556,23 +1604,24 @@ public:
ss << "}\n\t";
}
//The residual of mod outLoopSize
- for(unsigned int count=nCurWindowSize/outLoopSize*outLoopSize;
- count < nCurWindowSize; count++)
+ for (unsigned int count = nCurWindowSize / outLoopSize * outLoopSize;
+ count < nCurWindowSize; count++)
{
- ss << "i =" <<count<<";\n";
- if(count==nCurWindowSize/outLoopSize*outLoopSize){
+ ss << "i =" << count << ";\n";
+ if (count == nCurWindowSize / outLoopSize * outLoopSize)
+ {
for (unsigned i = 0; i < vSubArguments.size(); i++)
{
tmpCur = vSubArguments[i]->GetFormulaToken();
- if(ocPush==tmpCur->GetOpCode())
+ if (ocPush == tmpCur->GetOpCode())
{
- pCurDVR= static_cast<
- const formula::DoubleVectorRefToken *>(tmpCur);
- if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
+ pCurDVR = static_cast<
+ const
formula::DoubleVectorRefToken*>(tmpCur);
+ if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
{
temp4 << " currentCount";
temp4 << i;
- temp4 <<" =i+gid0+1;\n";
+ temp4 << " =i+gid0+1;\n";
}
else
{
@@ -1588,33 +1637,33 @@ public:
{
if (i)
temp4 << "*";
-
if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
+ if (ocPush ==
vSubArguments[i]->GetFormulaToken()->GetOpCode())
{
- temp4 <<"(";
- temp4 <<"(currentCount";
+ temp4 << "(";
+ temp4 << "(currentCount";
temp4 << i;
temp4 << ">";
- if(vSubArguments[i]->GetFormulaToken()->GetType() ==
+ if (vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svSingleVectorRef)
{
const formula::SingleVectorRefToken* pSVR =
- static_cast< const
formula::SingleVectorRefToken*>
+ static_cast<const
formula::SingleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken());
- temp4<<pSVR->GetArrayLength();
- temp4 << ")||isNan("<<vSubArguments[i]
+ temp4 << pSVR->GetArrayLength();
+ temp4 << ")||isNan(" << vSubArguments[i]
->GenSlidingWindowDeclRef();
temp4 << ")?0:";
temp4 <<
vSubArguments[i]->GenSlidingWindowDeclRef();
temp4 << ")";
}
- else if(vSubArguments[i]->GetFormulaToken()->GetType()
==
+ else if
(vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svDoubleVectorRef)
{
const formula::DoubleVectorRefToken* pSVR =
- static_cast< const
formula::DoubleVectorRefToken*>
+ static_cast<const
formula::DoubleVectorRefToken*>
(vSubArguments[i]->GetFormulaToken());
- temp4<<pSVR->GetArrayLength();
- temp4 << ")||isNan("<<vSubArguments[i]
+ temp4 << pSVR->GetArrayLength();
+ temp4 << ")||isNan(" << vSubArguments[i]
->GenSlidingWindowDeclRef(true);
temp4 << ")?0:";
temp4 <<
vSubArguments[i]->GenSlidingWindowDeclRef(true);
@@ -1642,188 +1691,206 @@ public:
};
/// operator traits
-class OpNop: public Reduction {
+class OpNop : public Reduction
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& )
const SAL_OVERRIDE
{
return lhs;
}
- virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "nop"; }
+ virtual std::string BinFuncName() const SAL_OVERRIDE { return "nop"; }
};
-class OpCount: public Reduction {
+class OpCount : public Reduction
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &rhs)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& rhs )
const SAL_OVERRIDE
{
std::stringstream ss;
- ss << "(isNan(" << lhs << ")?"<<rhs<<":"<<rhs<<"+1.0)";
+ ss << "(isNan(" << lhs << ")?" << rhs << ":" << rhs << "+1.0)";
return ss.str();
}
- virtual std::string BinFuncName(void) const SAL_OVERRIDE { return
"fcount"; }
+ virtual std::string BinFuncName() const SAL_OVERRIDE { return "fcount"; }
};
-class OpEqual: public Binary {
+class OpEqual : public Binary
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &rhs)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& rhs )
const SAL_OVERRIDE
{
std::stringstream ss;
- ss << "strequal("<< lhs << "," << rhs <<")";
+ ss << "strequal(" << lhs << "," << rhs << ")";
return ss.str();
}
- virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "eq"; }
+ virtual std::string BinFuncName() const SAL_OVERRIDE { return "eq"; }
};
-class OpLessEqual: public Binary {
+class OpLessEqual : public Binary
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &rhs)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& rhs )
const SAL_OVERRIDE
{
std::stringstream ss;
- ss << "("<< lhs << "<=" << rhs <<")";
+ ss << "(" << lhs << "<=" << rhs << ")";
return ss.str();
}
- virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "leq"; }
+ virtual std::string BinFuncName() const SAL_OVERRIDE { return "leq"; }
};
-class OpLess: public Binary {
+
+class OpLess : public Binary
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &rhs)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& rhs )
const SAL_OVERRIDE
{
std::stringstream ss;
- ss << "("<< lhs << "<" << rhs <<")";
+ ss << "(" << lhs << "<" << rhs << ")";
return ss.str();
}
- virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "less"; }
+ virtual std::string BinFuncName() const SAL_OVERRIDE { return "less"; }
};
-class OpGreater: public Binary {
+class OpGreater : public Binary
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &rhs)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& rhs )
const SAL_OVERRIDE
{
std::stringstream ss;
- ss << "("<< lhs << ">" << rhs <<")";
+ ss << "(" << lhs << ">" << rhs << ")";
return ss.str();
}
- virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "gt"; }
+ virtual std::string BinFuncName() const SAL_OVERRIDE { return "gt"; }
};
-class OpSum: public Reduction {
+class OpSum : public Reduction
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &rhs)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& rhs )
const SAL_OVERRIDE
{
std::stringstream ss;
- ss << "((" << lhs <<")+("<< rhs<<"))";
+ ss << "((" << lhs << ")+(" << rhs << "))";
return ss.str();
}
- virtual std::string BinFuncName(void) const SAL_OVERRIDE { return "fsum"; }
+ virtual std::string BinFuncName() const SAL_OVERRIDE { return "fsum"; }
};
-class OpAverage: public Reduction {
+class OpAverage : public Reduction
+{
public:
- virtual std::string GetBottom(void) SAL_OVERRIDE { return "0"; }
- virtual std::string Gen2(const std::string &lhs, const std::string &rhs)
const SAL_OVERRIDE
+ virtual std::string GetBottom() SAL_OVERRIDE { return "0"; }
+ virtual std::string Gen2( const std::string& lhs, const std::string& rhs )
const SAL_OVERRIDE
{
std::stringstream ss;
- ss << "fsum_count(" << lhs <<","<< rhs<<", &nCount)";
+ ss << "fsum_count(" << lhs << "," << rhs << ", &nCount)";
return ss.str();
... etc. - the rest is truncated
_______________________________________________
Libreoffice-commits mailing list
libreoffice-comm...@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/libreoffice-commits
